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

[Durgadoss R]

================
@@ -6247,6 +6176,98 @@ def NVVM_Tcgen05MMAWsSparseOp : NVVM_Op<"tcgen05.mma.ws.sp",
   }];
 }
 
+def SIMTFloatType : AnyTypeOf<[F16, BF16, F32, F64, VectorOfLengthAndType<[2], [F16, BF16, F32, F64]>]>;
+
+class ResultAtLeastAsWideAs<string operandArg> :
+  TypesMatchWith<"result type must be at least as wide as " # operandArg # " operand",
+    operandArg, "res", "$_self",
+    "::mlir::NVVM::isResultTypeAtLeastAsWideAsOperand">;
+
+class AllScalarsOrAllVectors<string lhsArg, string rhsArg> :
+  TypesMatchWith<"cannot mix vector and scalar operands", 
+    lhsArg, rhsArg, "$_self",
+    "::mlir::NVVM::areBothScalarsOrBothVectors">;
+
+class NVVM_FloatBinaryOp<string mnemonic, list<Trait> traits = []> :
+    NVVM_Op<mnemonic, traits # [Pure, ResultsAreFloatLike,
+                                ResultAtLeastAsWideAs<"lhs">,
+                                ResultAtLeastAsWideAs<"rhs">,
+                                AllScalarsOrAllVectors<"lhs", "rhs">]>,
+    Arguments<(ins SIMTFloatType:$lhs, SIMTFloatType:$rhs,
+      DefaultValuedAttr<FPRoundingModeAttr, "FPRoundingMode::NONE">:$rnd,
+      DefaultValuedAttr<SaturationModeAttr, "SaturationMode::NONE">:$sat,
+      DefaultValuedAttr<BoolAttr, "false">:$ftz)>,
+    Results<(outs SIMTFloatType:$res)> {
+  let assemblyFormat = "$lhs `,` $rhs attr-dict `:` type(operands) `->` type($res)";
+}
+
+def NVVM_AddFOp : NVVM_FloatBinaryOp<"addf", [Commutative]> {
+  let summary = [{
+    Performs floating point addition operation with support for mixed precision 
+    operands
+  }];
+  let description = [{
+    The `nvvm.addf` operation performs floating point addition of two operands.
+
+    The rounding mode is specified by the `rnd` attribute, saturation mode by 
+    the `sat` attribute, and flush-to-zero by the `ftz` attribute.
+
+    **Type constraints:**
+    - The result type must be at least as wide as both operands.
+    - Operands and result must be all scalars or all vectors (no mixing).
+    - When operands are narrower than the result, they are extended to the 
+      result type before addition. When this occurs, the modifiers that are 
+      supported for the addition will depend upon the result type.
----------------
durga4github wrote:

"...supported for the addition is determined by the result type."

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