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

[Guray Ozen]

================
@@ -3072,6 +3072,91 @@ LogicalResult NVVM::TensormapReplaceOp::verify() {
   return success();
 }
 
+LogicalResult NVVM::FAddOp::verify() {
+  mlir::Type resFType = getRes().getType();
+  mlir::Type lhsFType = getLhs().getType();
+  mlir::Type rhsFType = getRhs().getType();
+  mlir::NVVM::FPRoundingMode rndMode = getRnd();
+  mlir::NVVM::SaturationMode satMode = getSat();
+  bool isFTZ = getFtz();
+
+  auto getBaseFType = [](Type type) -> Type {
+    if (isa<VectorType>(type))
+      return cast<VectorType>(type).getElementType();
+    return type;
+  };
+
+  auto resBaseFType = getBaseFType(resFType);
+  auto lhsBaseFType = getBaseFType(lhsFType);
+  auto rhsBaseFType = getBaseFType(rhsFType);
+
+  // Supported operand types based on result types are:
+  // Result Type    : Operand Type(s)
+  // f64            : f64, f32, f16, bf16
+  // f32            : f32, f16, bf16
+  // f16            : f16
+  // bf16           : bf16
+  // vector<2xf64>  : vector<2x{f64, f32, f16, bf16}
+  // vector<2xf32>  : vector<2x{f32, f16, bf16}
+  // vector<2xf16>  : vector<2xf16>
+  // vector<2xbf16> : vector<2xbf16>
+
+  bool sameTypeOperation =
+      llvm::all_equal({lhsBaseFType, rhsBaseFType, resBaseFType});
+
+  if (!llvm::all_equal({isa<VectorType>(resFType), isa<VectorType>(lhsFType),
+                        isa<VectorType>(rhsFType)}))
+    return emitOpError("cannot mix vector and scalar types for floating point "
+                       "addition operation");
+
+  if (resBaseFType.getIntOrFloatBitWidth() <
+      std::max(lhsBaseFType.getIntOrFloatBitWidth(),
+               rhsBaseFType.getIntOrFloatBitWidth()))
+    return emitOpError("result type must be at least as wide as the operands");
----------------
grypp wrote:

If we have a tablegen class, we could move this kind checks there.  

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