[Mlir-commits] [mlir] [MLIR][NVVM] Add nvvm.fma Op (PR #184776)

[Durgadoss R]

================
@@ -3099,6 +3099,56 @@ LogicalResult NVVM::AddFOp::verify() {
   return success();
 }
 
+LogicalResult NVVM::FmaOp::verify() {
+  auto opType = getRes().getType();
+  mlir::NVVM::FPRoundingMode rndMode = getRnd();
+  mlir::NVVM::SaturationMode satMode = getSat();
+  bool isFTZ = getFtz();
+  bool isRelu = getRelu();
+  bool hasOOB = getOob();
+
+  auto getBaseFType = [](Type type) -> Type {
+    if (isa<VectorType>(type))
+      return cast<VectorType>(type).getElementType();
+    return type;
+  };
+
+  auto opBaseType = getBaseFType(opType);
+
+  if (rndMode == NVVM::FPRoundingMode::NONE)
+    return emitOpError("rounding mode must be specified");
+
+  if (isRelu && satMode == NVVM::SaturationMode::SAT)
+    return emitOpError("relu and saturation are not supported together");
+
+  if (hasOOB && (satMode == NVVM::SaturationMode::SAT || isFTZ))
+    return emitOpError("oob is not supported with saturation or FTZ");
+
+  if (!(opBaseType.isF16() || opBaseType.isBF16()) && (isRelu || hasOOB))
+    return emitOpError("relu and oob are only supported for f16 and bf16 fused "
+                       "multiply-add operations");
----------------
durga4github wrote:

Since it is emitOpError, (i..e that will tell us which Op it is), we don't need to suffix "fused mul-add ops".
The error-strings will be shorter

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