[Mlir-commits] [mlir] [MLIR][NVVM] Add nvvm.fadd and nvvm.fsub Ops (PR #179162)

[Durgadoss R]

================
@@ -6247,6 +6155,77 @@ def NVVM_Tcgen05MMAWsSparseOp : NVVM_Op<"tcgen05.mma.ws.sp",
   }];
 }
 
+def NVVM_FloatAdditionOp : 
+  NVVM_SingleResultIntrinsicOp<"fadd", [Pure, Commutative]> {
+  let summary = [{
+    Performs floating point addition operation with support for mixed precision 
+    operands
+  }];
+  let description = [{
+    The `nvvm.fadd` operation performs floating point addition of two operands.
+
+    The rounding mode to be used is specified by the `rnd` attribute, 
+    saturation mode by the `sat` attribute, and FTZ by the `ftz` unit attribute.
+    
+    The result type must be at least as wide as the operands. The operands are 
+    converted to the result type before addition if it is wider.
+    
+    For more information, see PTX ISA - [floating point addition](https://docs.nvidia.com/cuda/parallel-thread-execution/#floating-point-instructions-add), 
+    [half-precision floating point addition](https://docs.nvidia.com/cuda/parallel-thread-execution/#half-precision-floating-point-instructions-add), 
+    [mixed precision floating point addition](https://docs.nvidia.com/cuda/parallel-thread-execution/#mixed-precision-floating-point-instructions-add).
+  }];
+  let arguments = (ins
+    AnyTypeOf<[F16, BF16, F32, F64, VectorOfLengthAndType<[2], [F16, BF16]>]>:$lhs,
+    AnyTypeOf<[F16, BF16, F32, F64, VectorOfLengthAndType<[2], [F16, BF16]>]>:$rhs,
+    DefaultValuedAttr<FPRoundingModeAttr, "FPRoundingMode::NONE">:$rnd,
+    DefaultValuedAttr<SaturationModeAttr, "SaturationMode::NONE">:$sat,
+    UnitAttr:$ftz
+  );
+  let results = (outs AnyTypeOf<[F16, BF16, F32, F64, VectorOfLengthAndType<[2], [F16, BF16]>]>:$res);
+  let assemblyFormat = "$lhs `,` $rhs attr-dict `:` type(operands) `->` type($res)";
+  let hasVerifier = 1;
----------------
durga4github wrote:

optional for us to think through:
Coming back here after reviewing all the test files:
Now, I wonder whether doing the high-level Op is the correct thing here..
It seems there are many invalid cases in result/arg mixed types.

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