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

[Guray Ozen]

================
@@ -6247,6 +6176,80 @@ def NVVM_Tcgen05MMAWsSparseOp : NVVM_Op<"tcgen05.mma.ws.sp",
   }];
 }
 
+def NVVMFloatType : AnyTypeOf<[F16, BF16, F32, F64, VectorOfLengthAndType<[2], [F16, BF16]>]>;
+
+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. When the type of 
+    the `res` is wider than the type of the operands, the operands are first 
+    converted to the result type, and then the addition is performed.
+    
+    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
+    NVVMFloatType:$lhs,
+    NVVMFloatType:$rhs,
+    DefaultValuedAttr<FPRoundingModeAttr, "FPRoundingMode::NONE">:$rnd,
+    DefaultValuedAttr<SaturationModeAttr, "SaturationMode::NONE">:$sat,
+    UnitAttr:$ftz
+  );
+  let results = (outs NVVMFloatType:$res);
+  let assemblyFormat = "$lhs `,` $rhs attr-dict `:` type(operands) `->` type($res)";
+  let hasVerifier = 1;
+  
+  let llvmBuilder = [{
+    auto [ID, args] = NVVM::FloatAdditionOp::getIntrinsicIDAndArgs(*op, moduleTranslation, builder);
+    if(ID != llvm::Intrinsic::not_intrinsic) {
+      llvm::Value *addResult = createIntrinsicCall(builder, ID, args);
+      $res = ($_resultType->getScalarSizeInBits() > 
+              addResult->getType()->getScalarSizeInBits()) 
+             ? builder.CreateFPExt(addResult, $_resultType) : addResult;
+    }
+  }];
+}
+
+def NVVM_FloatSubtractionOp : 
----------------
grypp wrote:

nit : let's keep name of the op and calss same 

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