[llvm] [mlir] [NVVM][NVPTX] Change TMA Tensor reduction ops to use flag for reduction ops (PR #213638)

[Durgadoss R via llvm-commits]

================
@@ -2062,6 +2062,75 @@ described in the `s2g.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.reduce.tile.[1-5]d`'
+
+##### Syntax:
+
+```llvm
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i32 %red_op, i1 %flag_ch)
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.2d(..., i32 %d0, i32 %d1, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.3d(..., i32 %d0, i32 %d1, i32 %d2, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)
+```
+
+##### Overview:
+
+The '`@llvm.nvvm.cp.async.bulk.tensor.reduce.tile.[1-5]d`' intrinsics
+correspond to the `cp.reduce.async.bulk.tensor.[1-5]d.global.shared::cta.*`
+set of PTX instructions. These instructions initiate an asynchronous reduction
+operation of tensor data in global memory with the tensor data in shared::cta
+memory, using `tile` mode. The dimension of the tensor data ranges from 1d to
+5d with the coordinates specified by the `i32 %d0 ... i32 %d4` arguments. The
+`i32 %red_op` argument selects the reduction operation to perform. It must be
+a compile-time constant in the half-open range `[0, 8)`, with the following
+encoding:
+
+| `red_op` | Reduction |
+|---------:|-----------|
+| 0 | add |
+| 1 | min |
+| 2 | max |
+| 3 | inc |
+| 4 | dec |
+| 5 | and |
+| 6 | or |
----------------
durga4github wrote:

nit: possible to align the spacing with the line above?

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