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

[Rajat Bajpai 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
----------------
rajatbajpai wrote:

Done

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