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

[Durgadoss R]

================
@@ -1795,16 +1795,71 @@ instructions. These instructions initiate an asynchronous copy of bulk data from
 global memory to shared::cluster memory. The 32-bit operand `%size` specifies
 the amount of memory to be copied and it must be a multiple of 16.
 
-- The last two arguments to these intrinsics are boolean flags indicating
-  support for cache_hint and/or multicast modifiers. These flag arguments must
-  be compile-time constants. The backend looks through these flags and lowers
-  the intrinsics appropriately.
-- The Nth argument (denoted by `i1 %flag_ch`) when set, indicates a valid
-  cache_hint (`i64 %ch`) and generates the `.L2::cache_hint` variant of the
-  PTX instruction.
-- The [N-1]th argument (denoted by `i1 %flag_mc`) when set, indicates the
-  presence of a multicast mask (`i16 %mc`) and generates the PTX instruction
-  with the `.multicast::cluster` modifier.
+- The trailing `%flag_mc`, `%flag_ch`, and `%validate_pattern` arguments control
+  the multicast and cache-hint modifiers and the data-validity reporting
+  pattern. They must be compile-time constants.
+- The argument denoted by `i1 %flag_ch`, when set, indicates a valid cache_hint
+  (`i64 %ch`) and generates the `.L2::cache_hint` variant of the PTX
+  instruction.
+- The argument denoted by `i1 %flag_mc`, when set, indicates the presence of a
+  multicast mask (`%mc`) and generates the PTX instruction with the
+  `.multicast::cluster` modifier. An `i16` mask selects the default/16-bit
+  multicast form while an `i32` mask selects the `.multicast::cluster::32b`
+  variant. Only `i16` and `i32` mask types are supported. The 32-bit form
+  requires `sm_107f` or later in the same family and is supported only on PTX
+  9.4 or later.
+- The trailing `i32 %validate_pattern` flag selects the data-validity reporting
+  mechanism; see
+  [TMA Global-to-Shared Data-Validity Reporting](#tma-global-to-shared-data-validity-reporting).
+
+For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk).
+
+#### '`llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed`'
+
+##### Syntax:
+
+```llvm
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i16(ptr addrspace(7) %dst, ptr addrspace(3) %mbar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %scope, i32 %validate_pattern)
+declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster.relaxed.i32(..., i32 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %scope, i32 %validate_pattern)
----------------
durga4github wrote:

flag_scope and flag_valid_pattern to make the naming consistent with other flag_ operand docs.
(and a few places below..)

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