[llvm] [mlir] [LLVM][NVPTX] Add Rubin extensions for G2S Tensor intrinsics (PR #220029)

[Durgadoss R via llvm-commits]

================
@@ -2369,13 +2402,167 @@ tensor operation. For the `im2col.w` and `im2col.w.128` mode, the number of
 offsets is always 2, denoted by `i16 %wHalo` and `i16 %wOffset` arguments.
 For more information on `im2col.w` and `im2col.w.128` modes, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-im2col-w-w128-modes).
 
-- The last argument to these intrinsics is a boolean flag indicating support
-  for cache_hint. This flag argument must be a compile-time constant. When set,
-  it indicates a valid cache_hint (`i64 %ch`) and generates the
-  `.L2::cache_hint` variant of the PTX instruction.
+The last two operands are a boolean cache-hint flag and the data-validity
+reporting pattern, with the same functionality as described in the `tile`
+mode intrinsics above.
 
 For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor).
 
+#### '`llvm.nvvm.cp.async.bulk.tensor.g2s.tile*.override*[1-5]d`' Intrinsics
+
+##### Syntax:
+
+```llvm
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.override.addr.1d.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr %tensor_map, ptr addrspace(1) %override_addr, i32 %d0, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %cta_group, i32 %validate_pattern)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.override.addr.1d.i32(..., i32 %mc, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.override.addr.2d.{i16,i32}(..., i32 %d0, i32 %d1, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.override.addr.3d.{i16,i32}(..., i32 %d0, i32 %d1, i32 %d2, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.override.addr.4d.{i16,i32}(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.override.addr.5d.{i16,i32}(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)
+
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.gather4.override.addr.2d.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr %tensor_map, ptr addrspace(1) %override_addr, i32 %x0, i32 %y0, i32 %y1, i32 %y2, i32 %y3, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %cta_group, i32 %validate_pattern)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.gather4.override.addr.2d.i32(..., i32 %mc, ...)
+
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.override.addr.dim.1d.i32(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr %tensor_map, ptr addrspace(1) %override_addr, i16 %ts0, i32 %d0, i32 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %cta_group, i32 %validate_pattern)
+
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.override.addr.dim.stride.2d.i32(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr %tensor_map, ptr addrspace(1) %override_addr, i16 %ts0, i16 %ts1, i32 %stride0, i16 %upper_stride, i32 %d0, i32 %d1, i32 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %cta_group, i32 %validate_pattern)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.override.addr.dim.stride.3d.i32(..., i16 %ts0, i16 %ts1, i16 %ts2, i32 %stride0, i32 %stride1, i16 %upper_stride, i32 %d0, i32 %d1, i32 %d2, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.override.addr.dim.stride.4d.i32(..., i16 %ts0, i16 %ts1, i16 %ts2, i16 %ts3, i32 %stride0, i32 %stride1, i32 %stride2, i16 %upper_stride, i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.override.addr.dim.stride.5d.i32(..., i16 %ts0, i16 %ts1, i16 %ts2, i16 %ts3, i16 %ts4, i32 %stride0, i32 %stride1, i32 %stride2, i32 %stride3, i16 %upper_stride, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)
+```
+
+##### Overview:
+
+These intrinsic asynchronously copy tile-mode tensor data from Global memory to Shared Cluster memory
+while allowing overriding properties in the opaque tensor map with
+explicit operands. The override operands and their restrictions are described
+in [Tensor-Map Property Overrides](#tensor-map-property-overrides).
+
+- `%flag_mc` controls whether the multicast mask (i16/i32 `%mc`) is emitted.
+- `%flag_ch` controls whether `%ch` cache hint is emitted.
+- `%cta_group` selects CTA-group qualifier, it can represent three values:
+    | Value | CTA Group Modifier              |
+    |-------|---------------------------------|
+    | 0     | no cta group modifier (default) |
+    | 1     | `cta_group::1`                  |
+    | 2     | `cta_group::2`                  |
+- The last `i32 %validate_pattern` flag is described in
+[TMA Global-to-Shared Data-Validity Reporting](#tma-global-to-shared-data-validity-reporting).
+
+All above described flag operands must be compile time constants.
+
+These intrinsics lower to the `.override::global_address` form, optionally
+with `.override::global_dim` or `.override::global_dim_stride`, and require
+`sm_107f` and PTX 9.4 or later.
+
+#### '`llvm.nvvm.cp.async.bulk.tensor.g2s.im2col*.override.addr.[3-5]d`' Intrinsics
+
+##### Syntax:
+
+```llvm
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.override.addr.3d.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr %tensor_map, ptr addrspace(1) %override_addr, i32 %d0, i32 %d1, i32 %d2, i16 %im2col0, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %cta_group, i32 %validate_pattern)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.override.addr.3d.i32(..., i32 %mc, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.override.addr.4d.{i16,i32}(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 %im2col0, i16 %im2col1, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.override.addr.5d.{i16,i32}(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 %im2col0, i16 %im2col1, i16 %im2col2, ...)
+
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.w.override.addr.3d.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr %tensor_map, ptr addrspace(1) %override_addr, i32 %d0, i32 %d1, i32 %d2, i16 %wHalo, i16 %wOffset, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %cta_group, i32 %validate_pattern)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.w.override.addr.3d.i32(..., i32 %mc, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.w.override.addr.4d.{i16,i32}(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.w.override.addr.5d.{i16,i32}(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)
+
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.w.128.override.addr.3d.i16(ptr addrspace(7) %dst, ptr addrspace(3) %bar, ptr %tensor_map, ptr addrspace(1) %override_addr, i32 %d0, i32 %d1, i32 %d2, i16 %wHalo, i16 %wOffset, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch, i32 %cta_group, i32 %validate_pattern)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.w.128.override.addr.3d.i32(..., i32 %mc, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.w.128.override.addr.4d.{i16,i32}(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.w.128.override.addr.5d.{i16,i32}(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)
+```
+
+##### Overview:
+
+These intrinsic asynchronously copy im2col-mode tensor data from Global memory to Shared Cluster memory
----------------
durga4github wrote:

same here too

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