[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)";
+}
----------------
durga4github wrote:

This looks a lot more terse and clean now!

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