[Mlir-commits] [mlir] [mlir][NVVM] Add support for tcgen05.ld.red Op (PR #177330)

[Pradeep Kumar]

================
@@ -5322,6 +5322,111 @@ def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld", [NVVMRequiresSMa<[100, 101]>]> {
   }];
 }
 
+//===----------------------------------------------------------------------===//
+// NVVM tcgen05.ld.red Op
+//===----------------------------------------------------------------------===//
+
+def Tcgen05LdRedMin: I32EnumAttrCase<"MIN", 0, "min">;
+def Tcgen05LdRedMax: I32EnumAttrCase<"MAX", 1, "max">;
+
+def Tcgen05LdRedOperation: I32EnumAttr<
+  "Tcgen05LdRedOperation",
+  "tcgen05.ld.red reduction operation",
+  [Tcgen05LdRedMin, Tcgen05LdRedMax]> {
+  let cppNamespace = "::mlir::NVVM";
+  let genSpecializedAttr = 0;
+}
+
+def Tcgen05LdRedOperationAttr:
+      EnumAttr<NVVM_Dialect, Tcgen05LdRedOperation, "tcgen05_ld_red_op"> {
+  let assemblyFormat = "`<` $value `>`";
+}
+
+def NVVM_Tcgen05LdRedOp : NVVM_Op<"tcgen05.ld.red", [NVVMRequiresSMa<[101]>]> {
+  let summary = "tensor memory load and reduce instructions";
+  let arguments = (ins
+    Tcgen05LdStShapeAttr:$shape,
+    Tcgen05LdRedOperationAttr:$op,
+    UnitAttr:$abs,
+    UnitAttr:$nan,
+    LLVM_PointerTensor:$addr,
+    Optional<I64>:$offset
+  );
+
+  let results = (outs VectorOfLengthAndType<[2, 4, 8, 16, 32, 64, 128],
+                                            [I32, F32]>:$data,
+                      AnyTypeOf<[I32, F32]>:$redVal);
+
+  let assemblyFormat = [{
+    $addr (`,` $offset^)? attr-dict `:` type($data) `,` type($redVal)
+  }];
+
+  let description = [{
+    Instruction `tcgen05.ld.red` asynchronously loads data from the Tensor
+    Memory at the location specified by the 32-bit address operand `addr` into
+    the destination register `data`, collectively across all threads of the
+    warps. The operaiton also performs reduction operation specified by `op` on
----------------
schwarzschild-radius wrote:

Thanks for catching that! Fixed the description in the latest revision

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