[Mlir-commits] [mlir] 95a9647 - [MLIR][NVVM] Add nvvm.fma Op (#184776)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Sun Mar 8 20:46:21 PDT 2026
Author: Srinivasa Ravi
Date: 2026-03-09T09:16:16+05:30
New Revision: 95a9647f2ac5827042a4755c8c022bcebaf2bca8
URL: https://github.com/llvm/llvm-project/commit/95a9647f2ac5827042a4755c8c022bcebaf2bca8
DIFF: https://github.com/llvm/llvm-project/commit/95a9647f2ac5827042a4755c8c022bcebaf2bca8.diff
LOG: [MLIR][NVVM] Add nvvm.fma Op (#184776)
Adds `nvvm.fma` Op to the NVVM dialect to perform fused multiply-add
operations.
PTX ISA Reference:
1.
https://docs.nvidia.com/cuda/parallel-thread-execution/#floating-point-instructions-fma
2.
https://docs.nvidia.com/cuda/parallel-thread-execution/#half-precision-floating-point-instructions-fma
Added:
mlir/test/Target/LLVMIR/nvvm/fma/fma.mlir
mlir/test/Target/LLVMIR/nvvm/fma/fma_invalid.mlir
mlir/test/Target/LLVMIR/nvvm/fma/fma_vector.mlir
Modified:
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
Removed:
################################################################################
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index f8e1ab38e80d4..57a0c67e82c47 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -6380,6 +6380,49 @@ def NVVM_SubFOp : NVVM_FloatBinaryOp<"subf"> {
}];
}
+def NVVM_FmaOp : NVVM_Op<"fma", [Pure, SameOperandsAndResultType]> {
+ let summary = [{
+ Performs floating point fused multiply-add operation with support for mixed
+ precision operands
+ }];
+ let description = [{
+ The `nvvm.fma` operation performs floating point fused multiply-add of
+ three operands of the same type.
+
+ The rounding mode is specified by the `rnd` attribute, saturation mode by
+ the `sat` attribute, flush-to-zero by the `ftz` attribute, and ReLU by the
+ `relu` attribute.
+
+ Out-of-bounds (OOB) behavior is controlled by the `oob` attribute. `oob`
+ clamps the result to 0 if either of the operands is `OOB NaN` (see [Tensors](https://docs.nvidia.com/cuda/parallel-thread-execution/#tensors)).
+
+ For more information, see PTX ISA:
+ - [floating point fused multiply-add](https://docs.nvidia.com/cuda/parallel-thread-execution/#floating-point-instructions-fma)
+ - [half-precision floating point fused multiply-add](https://docs.nvidia.com/cuda/parallel-thread-execution/#half-precision-floating-point-instructions-fma)
+ }];
+ let arguments = (ins
+ SIMTFloatType:$a,
+ SIMTFloatType:$b,
+ SIMTFloatType:$c,
+ FPArithRoundingMode:$rnd,
+ DefaultValuedAttr<SaturationModeSatOrNone, "SaturationMode::NONE">:$sat,
+ DefaultValuedAttr<BoolAttr, "false">:$ftz,
+ DefaultValuedAttr<BoolAttr, "false">:$relu,
+ DefaultValuedAttr<BoolAttr, "false">:$oob
+ );
+ let results = (outs SIMTFloatType:$res);
+ let assemblyFormat = "$a `,` $b `,` $c attr-dict `:` type($a)";
+ let hasVerifier = 1;
+
+ let extraClassDeclaration = [{
+ static void lowerFmaToLLVMIR(
+ Operation &op, LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder);
+ }];
+ let llvmBuilder = [{
+ NVVM::FmaOp::lowerFmaToLLVMIR(*op, moduleTranslation, builder);
+ }];
+}
+
//===----------------------------------------------------------------------===//
// NVVM tensormap.replace Op
//===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 6ccd59cec65bc..7d49aa3878ebe 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -3104,6 +3104,53 @@ LogicalResult NVVM::AddFOp::verify() { return verifyAddSubFOp<AddFOp>(*this); }
LogicalResult NVVM::SubFOp::verify() { return verifyAddSubFOp<SubFOp>(*this); }
+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");
+
+ if (opBaseType.isF64() && (satMode != NVVM::SaturationMode::NONE || isFTZ))
+ return emitOpError("FTZ and saturation are not supported for f64 type");
+
+ if (opBaseType.isF16() && rndMode != NVVM::FPRoundingMode::RN)
+ return emitOpError(
+ "only RN rounding mode is supported for f16 and vector<2xf16>");
+
+ if (opBaseType.isBF16()) {
+ if (rndMode != NVVM::FPRoundingMode::RN)
+ return emitOpError(
+ "only RN rounding mode is supported for bf16 and vector<2xbf16>");
+ if (satMode != NVVM::SaturationMode::NONE || isFTZ)
+ return emitOpError(
+ "FTZ and saturation are not supported for bf16 and vector<2xbf16>");
+ }
+
+ return success();
+}
+
/// Packs the given `field` into the `result`.
/// The `result` is 64-bits and each `field` can be 32-bits or narrower.
static llvm::Value *
diff --git a/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
index 092643f408ce6..5e5f6700c9fd7 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp
@@ -446,6 +446,33 @@ getFenceProxySyncRestrictID(NVVM::MemOrderKind order) {
nvvm_fence_proxy_async_generic_release_sync_restrict_space_cta_scope_cluster;
}
+// Calls an LLVM intrinsic on the given operands. For f32/f64 vector types,
+// the intrinsic is called per-element and the results are packed back into a
+// vector. If retType is non-null, it is forwarded as the return-type
+// overload to `createIntrinsicCall`.
+static llvm::Value *
+createScalarizedIntrinsicCall(llvm::IRBuilderBase &builder,
+ llvm::Intrinsic::ID IID, llvm::Type *opTypeLLVM,
+ ArrayRef<llvm::Value *> operands,
+ llvm::Type *retType) {
+ if (opTypeLLVM->isVectorTy() && (opTypeLLVM->getScalarType()->isFloatTy() ||
+ opTypeLLVM->getScalarType()->isDoubleTy())) {
+ llvm::Value *result = llvm::PoisonValue::get(
+ llvm::FixedVectorType::get(opTypeLLVM->getScalarType(), 2));
+ for (int64_t i = 0; i < 2; ++i) {
+ llvm::SmallVector<llvm::Value *> scalarArgs;
+ for (llvm::Value *op : operands)
+ scalarArgs.push_back(
+ builder.CreateExtractElement(op, builder.getInt32(i)));
+ llvm::Value *res = createIntrinsicCall(builder, IID, retType, scalarArgs);
+ result = builder.CreateInsertElement(result, res, builder.getInt32(i));
+ }
+ return result;
+ }
+
+ return createIntrinsicCall(builder, IID, retType, operands);
+}
+
void NVVM::AddFOp::lowerAddFToLLVMIR(llvm::Value *argLHS, llvm::Value *argRHS,
Value res, NVVM::FPRoundingMode rndMode,
NVVM::SaturationMode satMode, bool isFTZ,
@@ -493,31 +520,9 @@ void NVVM::AddFOp::lowerAddFToLLVMIR(llvm::Value *argLHS, llvm::Value *argRHS,
llvm::Intrinsic::nvvm_add_rp_d, llvm::Intrinsic::nvvm_add_rz_d};
auto addIntrinsic = [&](llvm::Intrinsic::ID IID) -> llvm::Value * {
- auto createAddIntrinsicCall = [&](llvm::Intrinsic::ID IID, llvm::Value *LHS,
- llvm::Value *RHS) -> llvm::CallInst * {
- llvm::SmallVector<llvm::Value *, 2> callArgs;
- callArgs.push_back(LHS);
- callArgs.push_back(RHS);
- return createIntrinsicCall(builder, IID, callArgs);
- };
-
- if (isVectorOp && (opTypeLLVM->getScalarType()->isFloatTy() ||
- opTypeLLVM->getScalarType()->isDoubleTy())) {
- llvm::Value *result = llvm::PoisonValue::get(
- llvm::FixedVectorType::get(opTypeLLVM->getScalarType(), 2));
- for (int64_t i = 0; i < 2; ++i) {
- llvm::Value *lhsElemi =
- builder.CreateExtractElement(argLHS, builder.getInt32(i));
- llvm::Value *rhsElemi =
- builder.CreateExtractElement(argRHS, builder.getInt32(i));
- llvm::Value *sum = createAddIntrinsicCall(IID, lhsElemi, rhsElemi);
- result = builder.CreateInsertElement(result, sum, builder.getInt32(i));
- };
- return result;
- }
-
- return createAddIntrinsicCall(IID, argLHS, argRHS);
- }; // addIntrinsic end
+ return createScalarizedIntrinsicCall(builder, IID, opTypeLLVM,
+ {argLHS, argRHS}, opTypeLLVM);
+ };
// f16 + f16 -> f16 / vector<2xf16> + vector<2xf16> -> vector<2xf16>
// FIXME: Allow lowering to add.rn.ftz.f16x2 and add.rn.ftz.f16 here when the
@@ -557,6 +562,122 @@ void NVVM::AddFOp::lowerAddFToLLVMIR(llvm::Value *argLHS, llvm::Value *argRHS,
}
}
+void NVVM::FmaOp::lowerFmaToLLVMIR(Operation &op, LLVM::ModuleTranslation &mt,
+ llvm::IRBuilderBase &builder) {
+ auto thisOp = cast<NVVM::FmaOp>(op);
+ mlir::NVVM::FPRoundingMode rndMode = thisOp.getRnd();
+ unsigned rndIndex = static_cast<unsigned>(rndMode) - 1; // 1-4 mapped to 0-3
+ mlir::NVVM::SaturationMode satMode = thisOp.getSat();
+ bool isFTZ = thisOp.getFtz();
+ bool isRelu = thisOp.getRelu();
+ bool isSat = satMode == NVVM::SaturationMode::SAT;
+ bool isOOB = thisOp.getOob();
+
+ mlir::Type opType = thisOp.getRes().getType();
+ llvm::Type *opTypeLLVM = mt.convertType(opType);
+ bool isVectorFma = opTypeLLVM->isVectorTy();
+
+ llvm::Value *argA = mt.lookupValue(thisOp.getA());
+ llvm::Value *argB = mt.lookupValue(thisOp.getB());
+ llvm::Value *argC = mt.lookupValue(thisOp.getC());
+
+ static constexpr llvm::Intrinsic::ID f16IDs[] = {
+ llvm::Intrinsic::nvvm_fma_rn_f16,
+ llvm::Intrinsic::nvvm_fma_rn_f16x2,
+ llvm::Intrinsic::nvvm_fma_rn_ftz_f16,
+ llvm::Intrinsic::nvvm_fma_rn_ftz_f16x2,
+ llvm::Intrinsic::nvvm_fma_rn_sat_f16,
+ llvm::Intrinsic::nvvm_fma_rn_sat_f16x2,
+ llvm::Intrinsic::nvvm_fma_rn_ftz_sat_f16,
+ llvm::Intrinsic::nvvm_fma_rn_ftz_sat_f16x2,
+ llvm::Intrinsic::nvvm_fma_rn_relu_f16,
+ llvm::Intrinsic::nvvm_fma_rn_relu_f16x2,
+ llvm::Intrinsic::nvvm_fma_rn_ftz_relu_f16,
+ llvm::Intrinsic::nvvm_fma_rn_ftz_relu_f16x2};
+
+ static constexpr llvm::Intrinsic::ID bf16IDs[] = {
+ llvm::Intrinsic::nvvm_fma_rn_bf16, llvm::Intrinsic::nvvm_fma_rn_bf16x2,
+ llvm::Intrinsic::nvvm_fma_rn_relu_bf16,
+ llvm::Intrinsic::nvvm_fma_rn_relu_bf16x2};
+
+ static constexpr llvm::Intrinsic::ID f32IDs[] = {
+ llvm::Intrinsic::nvvm_fma_rn_f,
+ llvm::Intrinsic::nvvm_fma_rm_f,
+ llvm::Intrinsic::nvvm_fma_rp_f,
+ llvm::Intrinsic::nvvm_fma_rz_f,
+ llvm::Intrinsic::nvvm_fma_rn_sat_f,
+ llvm::Intrinsic::nvvm_fma_rm_sat_f,
+ llvm::Intrinsic::nvvm_fma_rp_sat_f,
+ llvm::Intrinsic::nvvm_fma_rz_sat_f,
+ llvm::Intrinsic::nvvm_fma_rn_ftz_f,
+ llvm::Intrinsic::nvvm_fma_rm_ftz_f,
+ llvm::Intrinsic::nvvm_fma_rp_ftz_f,
+ llvm::Intrinsic::nvvm_fma_rz_ftz_f,
+ llvm::Intrinsic::nvvm_fma_rn_ftz_sat_f,
+ llvm::Intrinsic::nvvm_fma_rm_ftz_sat_f,
+ llvm::Intrinsic::nvvm_fma_rp_ftz_sat_f,
+ llvm::Intrinsic::nvvm_fma_rz_ftz_sat_f,
+ };
+
+ static constexpr llvm::Intrinsic::ID f64IDs[] = {
+ llvm::Intrinsic::nvvm_fma_rn_d, llvm::Intrinsic::nvvm_fma_rm_d,
+ llvm::Intrinsic::nvvm_fma_rp_d, llvm::Intrinsic::nvvm_fma_rz_d};
+
+ auto fmaIntrinsic = [&](llvm::Intrinsic::ID IID,
+ llvm::Type *retType) -> llvm::Value * {
+ return createScalarizedIntrinsicCall(
+ builder, IID, opTypeLLVM, {argA, argB, argC}, /*retType=*/retType);
+ };
+
+ // f16 + f16 -> f16 / vector<2xf16> + vector<2xf16> -> vector<2xf16>
+ if (opTypeLLVM->getScalarType()->isHalfTy()) {
+ llvm::Value *result;
+ if (isOOB) {
+ result = fmaIntrinsic(isRelu ? llvm::Intrinsic::nvvm_fma_rn_oob_relu
+ : llvm::Intrinsic::nvvm_fma_rn_oob,
+ opTypeLLVM);
+ } else {
+ unsigned index =
+ (isRelu << 3) | (isSat << 2) | (isFTZ << 1) |
+ isVectorFma; // Op verifier ensures that this index is valid
+ result = fmaIntrinsic(f16IDs[index], opTypeLLVM);
+ }
+ mt.mapValue(thisOp.getRes(), result);
+ return;
+ }
+
+ // bf16 + bf16 -> bf16 / vector<2xbf16> + vector<2xbf16> -> vector<2xbf16>
+ if (opTypeLLVM->getScalarType()->isBFloatTy()) {
+ llvm::Value *result;
+ if (isOOB) {
+ result = fmaIntrinsic(isRelu ? llvm::Intrinsic::nvvm_fma_rn_oob_relu
+ : llvm::Intrinsic::nvvm_fma_rn_oob,
+ opTypeLLVM);
+ } else {
+ unsigned index = (isRelu << 1) | isVectorFma;
+ result = fmaIntrinsic(bf16IDs[index], opTypeLLVM);
+ }
+ mt.mapValue(thisOp.getRes(), result);
+ return;
+ }
+
+ // f64 + f64 -> f64 / vector<2xf64> + vector<2xf64> -> vector<2xf64>
+ if (opTypeLLVM->getScalarType()->isDoubleTy()) {
+ mt.mapValue(thisOp.getRes(),
+ fmaIntrinsic(f64IDs[rndIndex], opTypeLLVM->getScalarType()));
+ return;
+ }
+
+ // f32 + f32 -> f32 / vector<2xf32> + vector<2xf32> -> vector<2xf32>
+ const unsigned numRndModes = 4; // RN, RM, RP, RZ
+ if (opTypeLLVM->getScalarType()->isFloatTy()) {
+ unsigned index = ((isFTZ << 1) | isSat) * numRndModes + rndIndex;
+ mt.mapValue(thisOp.getRes(),
+ fmaIntrinsic(f32IDs[index], opTypeLLVM->getScalarType()));
+ return;
+ }
+}
+
namespace {
/// Implementation of the dialect interface that converts operations belonging
/// to the NVVM dialect to LLVM IR.
diff --git a/mlir/test/Target/LLVMIR/nvvm/fma/fma.mlir b/mlir/test/Target/LLVMIR/nvvm/fma/fma.mlir
new file mode 100644
index 0000000000000..236175daff21e
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/nvvm/fma/fma.mlir
@@ -0,0 +1,114 @@
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
+
+llvm.func @fma_f16(%a: f16, %b: f16, %c: f16) -> f16 {
+ // CHECK-LABEL: define half @fma_f16(half %0, half %1, half %2) {
+ // CHECK-NEXT: %4 = call half @llvm.nvvm.fma.rn.f16(half %0, half %1, half %2)
+ // CHECK-NEXT: %5 = call half @llvm.nvvm.fma.rn.ftz.f16(half %0, half %1, half %4)
+ // CHECK-NEXT: %6 = call half @llvm.nvvm.fma.rn.sat.f16(half %0, half %1, half %5)
+ // CHECK-NEXT: %7 = call half @llvm.nvvm.fma.rn.ftz.sat.f16(half %0, half %1, half %6)
+ // CHECK-NEXT: %8 = call half @llvm.nvvm.fma.rn.relu.f16(half %0, half %1, half %7)
+ // CHECK-NEXT: %9 = call half @llvm.nvvm.fma.rn.ftz.relu.f16(half %0, half %1, half %8)
+ // CHECK-NEXT: %10 = call half @llvm.nvvm.fma.rn.oob.f16(half %0, half %1, half %9)
+ // CHECK-NEXT: %11 = call half @llvm.nvvm.fma.rn.oob.relu.f16(half %0, half %1, half %10)
+ // CHECK-NEXT: ret half %11
+ // CHECK-NEXT: }
+ %f0 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rn>} : f16
+ %f1 = nvvm.fma %a, %b, %f0 {rnd = #nvvm.fp_rnd_mode<rn>, ftz = true} : f16
+ %f2 = nvvm.fma %a, %b, %f1 {rnd = #nvvm.fp_rnd_mode<rn>, sat = #nvvm.sat_mode<sat>} : f16
+ %f3 = nvvm.fma %a, %b, %f2 {rnd = #nvvm.fp_rnd_mode<rn>, sat = #nvvm.sat_mode<sat>, ftz = true} : f16
+ %f4 = nvvm.fma %a, %b, %f3 {rnd = #nvvm.fp_rnd_mode<rn>, relu = true} : f16
+ %f5 = nvvm.fma %a, %b, %f4 {rnd = #nvvm.fp_rnd_mode<rn>, relu = true, ftz = true} : f16
+ %f6 = nvvm.fma %a, %b, %f5 {rnd = #nvvm.fp_rnd_mode<rn>, oob = true} : f16
+ %f7 = nvvm.fma %a, %b, %f6 {rnd = #nvvm.fp_rnd_mode<rn>, oob = true, relu = true} : f16
+ llvm.return %f7 : f16
+}
+
+llvm.func @fma_bf16(%a: bf16, %b: bf16, %c: bf16) -> bf16 {
+ // CHECK-LABEL: define bfloat @fma_bf16(bfloat %0, bfloat %1, bfloat %2) {
+ // CHECK-NEXT: %4 = call bfloat @llvm.nvvm.fma.rn.bf16(bfloat %0, bfloat %1, bfloat %2)
+ // CHECK-NEXT: %5 = call bfloat @llvm.nvvm.fma.rn.relu.bf16(bfloat %0, bfloat %1, bfloat %4)
+ // CHECK-NEXT: %6 = call bfloat @llvm.nvvm.fma.rn.oob.bf16(bfloat %0, bfloat %1, bfloat %5)
+ // CHECK-NEXT: %7 = call bfloat @llvm.nvvm.fma.rn.oob.relu.bf16(bfloat %0, bfloat %1, bfloat %6)
+ // CHECK-NEXT: ret bfloat %7
+ // CHECK-NEXT: }
+ %f0 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rn>} : bf16
+ %f1 = nvvm.fma %a, %b, %f0 {rnd = #nvvm.fp_rnd_mode<rn>, relu = true} : bf16
+ %f2 = nvvm.fma %a, %b, %f1 {rnd = #nvvm.fp_rnd_mode<rn>, oob = true} : bf16
+ %f3 = nvvm.fma %a, %b, %f2 {rnd = #nvvm.fp_rnd_mode<rn>, oob = true, relu = true} : bf16
+ llvm.return %f3 : bf16
+}
+
+llvm.func @fma_f32_rn(%a: f32, %b: f32, %c: f32) -> f32 {
+ // CHECK-LABEL: define float @fma_f32_rn(float %0, float %1, float %2) {
+ // CHECK-NEXT: %4 = call float @llvm.nvvm.fma.rn.f(float %0, float %1, float %2)
+ // CHECK-NEXT: %5 = call float @llvm.nvvm.fma.rn.ftz.f(float %0, float %1, float %4)
+ // CHECK-NEXT: %6 = call float @llvm.nvvm.fma.rn.sat.f(float %0, float %1, float %5)
+ // CHECK-NEXT: %7 = call float @llvm.nvvm.fma.rn.ftz.sat.f(float %0, float %1, float %6)
+ // CHECK-NEXT: ret float %7
+ // CHECK-NEXT: }
+ %f0 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rn>} : f32
+ %f1 = nvvm.fma %a, %b, %f0 {rnd = #nvvm.fp_rnd_mode<rn>, ftz = true} : f32
+ %f2 = nvvm.fma %a, %b, %f1 {rnd = #nvvm.fp_rnd_mode<rn>, sat = #nvvm.sat_mode<sat>} : f32
+ %f3 = nvvm.fma %a, %b, %f2 {rnd = #nvvm.fp_rnd_mode<rn>, sat = #nvvm.sat_mode<sat>, ftz = true} : f32
+ llvm.return %f3 : f32
+}
+
+llvm.func @fma_f32_rm(%a: f32, %b: f32, %c: f32) -> f32 {
+ // CHECK-LABEL: define float @fma_f32_rm(float %0, float %1, float %2) {
+ // CHECK-NEXT: %4 = call float @llvm.nvvm.fma.rm.f(float %0, float %1, float %2)
+ // CHECK-NEXT: %5 = call float @llvm.nvvm.fma.rm.ftz.f(float %0, float %1, float %4)
+ // CHECK-NEXT: %6 = call float @llvm.nvvm.fma.rm.sat.f(float %0, float %1, float %5)
+ // CHECK-NEXT: %7 = call float @llvm.nvvm.fma.rm.ftz.sat.f(float %0, float %1, float %6)
+ // CHECK-NEXT: ret float %7
+ // CHECK-NEXT: }
+ %f0 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rm>} : f32
+ %f1 = nvvm.fma %a, %b, %f0 {rnd = #nvvm.fp_rnd_mode<rm>, ftz = true} : f32
+ %f2 = nvvm.fma %a, %b, %f1 {rnd = #nvvm.fp_rnd_mode<rm>, sat = #nvvm.sat_mode<sat>} : f32
+ %f3 = nvvm.fma %a, %b, %f2 {rnd = #nvvm.fp_rnd_mode<rm>, sat = #nvvm.sat_mode<sat>, ftz = true} : f32
+ llvm.return %f3 : f32
+}
+
+llvm.func @fma_f32_rp(%a: f32, %b: f32, %c: f32) -> f32 {
+ // CHECK-LABEL: define float @fma_f32_rp(float %0, float %1, float %2) {
+ // CHECK-NEXT: %4 = call float @llvm.nvvm.fma.rp.f(float %0, float %1, float %2)
+ // CHECK-NEXT: %5 = call float @llvm.nvvm.fma.rp.ftz.f(float %0, float %1, float %4)
+ // CHECK-NEXT: %6 = call float @llvm.nvvm.fma.rp.sat.f(float %0, float %1, float %5)
+ // CHECK-NEXT: %7 = call float @llvm.nvvm.fma.rp.ftz.sat.f(float %0, float %1, float %6)
+ // CHECK-NEXT: ret float %7
+ // CHECK-NEXT: }
+ %f0 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rp>} : f32
+ %f1 = nvvm.fma %a, %b, %f0 {rnd = #nvvm.fp_rnd_mode<rp>, ftz = true} : f32
+ %f2 = nvvm.fma %a, %b, %f1 {rnd = #nvvm.fp_rnd_mode<rp>, sat = #nvvm.sat_mode<sat>} : f32
+ %f3 = nvvm.fma %a, %b, %f2 {rnd = #nvvm.fp_rnd_mode<rp>, sat = #nvvm.sat_mode<sat>, ftz = true} : f32
+ llvm.return %f3 : f32
+}
+
+llvm.func @fma_f32_rz(%a: f32, %b: f32, %c: f32) -> f32 {
+ // CHECK-LABEL: define float @fma_f32_rz(float %0, float %1, float %2) {
+ // CHECK-NEXT: %4 = call float @llvm.nvvm.fma.rz.f(float %0, float %1, float %2)
+ // CHECK-NEXT: %5 = call float @llvm.nvvm.fma.rz.ftz.f(float %0, float %1, float %4)
+ // CHECK-NEXT: %6 = call float @llvm.nvvm.fma.rz.sat.f(float %0, float %1, float %5)
+ // CHECK-NEXT: %7 = call float @llvm.nvvm.fma.rz.ftz.sat.f(float %0, float %1, float %6)
+ // CHECK-NEXT: ret float %7
+ // CHECK-NEXT: }
+ %f0 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rz>} : f32
+ %f1 = nvvm.fma %a, %b, %f0 {rnd = #nvvm.fp_rnd_mode<rz>, ftz = true} : f32
+ %f2 = nvvm.fma %a, %b, %f1 {rnd = #nvvm.fp_rnd_mode<rz>, sat = #nvvm.sat_mode<sat>} : f32
+ %f3 = nvvm.fma %a, %b, %f2 {rnd = #nvvm.fp_rnd_mode<rz>, sat = #nvvm.sat_mode<sat>, ftz = true} : f32
+ llvm.return %f3 : f32
+}
+
+llvm.func @fma_f64(%a: f64, %b: f64, %c: f64) -> f64 {
+ // CHECK-LABEL: define double @fma_f64(double %0, double %1, double %2) {
+ // CHECK-NEXT: %4 = call double @llvm.nvvm.fma.rn.d(double %0, double %1, double %2)
+ // CHECK-NEXT: %5 = call double @llvm.nvvm.fma.rm.d(double %0, double %1, double %4)
+ // CHECK-NEXT: %6 = call double @llvm.nvvm.fma.rp.d(double %0, double %1, double %5)
+ // CHECK-NEXT: %7 = call double @llvm.nvvm.fma.rz.d(double %0, double %1, double %6)
+ // CHECK-NEXT: ret double %7
+ // CHECK-NEXT: }
+ %f0 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rn>} : f64
+ %f1 = nvvm.fma %a, %b, %f0 {rnd = #nvvm.fp_rnd_mode<rm>} : f64
+ %f2 = nvvm.fma %a, %b, %f1 {rnd = #nvvm.fp_rnd_mode<rp>} : f64
+ %f3 = nvvm.fma %a, %b, %f2 {rnd = #nvvm.fp_rnd_mode<rz>} : f64
+ llvm.return %f3 : f64
+}
diff --git a/mlir/test/Target/LLVMIR/nvvm/fma/fma_invalid.mlir b/mlir/test/Target/LLVMIR/nvvm/fma/fma_invalid.mlir
new file mode 100644
index 0000000000000..ea92b707b65de
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/nvvm/fma/fma_invalid.mlir
@@ -0,0 +1,89 @@
+// RUN: mlir-translate --mlir-to-llvmir --split-input-file --verify-diagnostics %s
+
+// -----
+
+llvm.func @fma_invalid_rnd_mode(%a : f16, %b : f16, %c : f16) -> f16 {
+ // expected-error at +1 {{rounding mode must be specified}}
+ %f1 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<none>} : f16
+ llvm.return %f1 : f16
+}
+
+// -----
+
+llvm.func @fma_invalid_sat_mode(%a : f16, %b : f16, %c : f16) -> f16 {
+ // expected-error at +1 {{attribute 'sat' failed to satisfy constraint: Describes the saturation mode whose value is one of {none, sat}}}
+ %f1 = nvvm.fma %a, %b, %c {sat = #nvvm.sat_mode<satfinite>, rnd = #nvvm.fp_rnd_mode<rn>} : f16
+ llvm.return %f1 : f16
+}
+
+// -----
+
+llvm.func @fma_invalid_relu_sat(%a : f16, %b : f16, %c : f16) -> f16 {
+ // expected-error at +1 {{relu and saturation are not supported together}}
+ %f1 = nvvm.fma %a, %b, %c {relu = true, sat = #nvvm.sat_mode<sat>, rnd = #nvvm.fp_rnd_mode<rn>} : f16
+ llvm.return %f1 : f16
+}
+
+// -----
+
+llvm.func @fma_invalid_oob_sat(%a : f16, %b : f16, %c : f16) -> f16 {
+ // expected-error at +1 {{oob is not supported with saturation or FTZ}}
+ %f1 = nvvm.fma %a, %b, %c {oob = true, sat = #nvvm.sat_mode<sat>, rnd = #nvvm.fp_rnd_mode<rn>} : f16
+ llvm.return %f1 : f16
+}
+
+// -----
+
+llvm.func @fma_invalid_oob_f64(%a : f64, %b : f64, %c : f64) -> f64 {
+ // expected-error at +1 {{relu and oob are only supported for f16 and bf16}}
+ %f1 = nvvm.fma %a, %b, %c {oob = true, rnd = #nvvm.fp_rnd_mode<rn>} : f64
+ llvm.return %f1 : f64
+}
+
+// -----
+
+llvm.func @fma_invalid_relu_oob(%a : f32, %b : f32, %c : f32) -> f32 {
+ // expected-error at +1 {{relu and oob are only supported for f16 and bf16}}
+ %f1 = nvvm.fma %a, %b, %c {relu = true, rnd = #nvvm.fp_rnd_mode<rn>} : f32
+ llvm.return %f1 : f32
+}
+
+// -----
+
+llvm.func @fma_invalid_ftz_sat_f64(%a : f64, %b : f64, %c : f64) -> f64 {
+ // expected-error at +1 {{FTZ and saturation are not supported for f64 type}}
+ %f1 = nvvm.fma %a, %b, %c {ftz = true, sat = #nvvm.sat_mode<sat>, rnd = #nvvm.fp_rnd_mode<rn>} : f64
+ llvm.return %f1 : f64
+}
+
+// -----
+
+llvm.func @fma_invalid_v2f16_rnd_mode(%a : vector<2xf16>, %b : vector<2xf16>, %c : vector<2xf16>) -> vector<2xf16> {
+ // expected-error at +1 {{only RN rounding mode is supported for f16 and vector<2xf16>}}
+ %f1 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rm>} : vector<2xf16>
+ llvm.return %f1 : vector<2xf16>
+}
+
+// -----
+
+llvm.func @fma_invalid_v2bf16_rnd_mode(%a : vector<2xbf16>, %b : vector<2xbf16>, %c : vector<2xbf16>) -> vector<2xbf16> {
+ // expected-error at +1 {{only RN rounding mode is supported for bf16 and vector<2xbf16>}}
+ %f1 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rm>} : vector<2xbf16>
+ llvm.return %f1 : vector<2xbf16>
+}
+
+// -----
+
+llvm.func @fma_invalid_ftz_v2bf16(%a : vector<2xbf16>, %b : vector<2xbf16>, %c : vector<2xbf16>) -> vector<2xbf16> {
+ // expected-error at +1 {{FTZ and saturation are not supported for bf16 and vector<2xbf16>}}
+ %f1 = nvvm.fma %a, %b, %c {ftz = true, rnd = #nvvm.fp_rnd_mode<rn>} : vector<2xbf16>
+ llvm.return %f1 : vector<2xbf16>
+}
+
+// -----
+
+llvm.func @fma_invalid_sat_v2bf16(%a : vector<2xbf16>, %b : vector<2xbf16>, %c : vector<2xbf16>) -> vector<2xbf16> {
+ // expected-error at +1 {{FTZ and saturation are not supported for bf16 and vector<2xbf16>}}
+ %f1 = nvvm.fma %a, %b, %c {sat = #nvvm.sat_mode<sat>, rnd = #nvvm.fp_rnd_mode<rn>} : vector<2xbf16>
+ llvm.return %f1 : vector<2xbf16>
+}
diff --git a/mlir/test/Target/LLVMIR/nvvm/fma/fma_vector.mlir b/mlir/test/Target/LLVMIR/nvvm/fma/fma_vector.mlir
new file mode 100644
index 0000000000000..020bdcfc27705
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/nvvm/fma/fma_vector.mlir
@@ -0,0 +1,294 @@
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
+
+llvm.func @fma_f16(%a: vector<2xf16>, %b: vector<2xf16>, %c: vector<2xf16>) -> vector<2xf16> {
+ // CHECK-LABEL: define <2 x half> @fma_f16(<2 x half> %0, <2 x half> %1, <2 x half> %2) {
+ // CHECK-NEXT: %4 = call <2 x half> @llvm.nvvm.fma.rn.f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %2)
+ // CHECK-NEXT: %5 = call <2 x half> @llvm.nvvm.fma.rn.ftz.f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %4)
+ // CHECK-NEXT: %6 = call <2 x half> @llvm.nvvm.fma.rn.sat.f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %5)
+ // CHECK-NEXT: %7 = call <2 x half> @llvm.nvvm.fma.rn.ftz.sat.f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %6)
+ // CHECK-NEXT: %8 = call <2 x half> @llvm.nvvm.fma.rn.relu.f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %7)
+ // CHECK-NEXT: %9 = call <2 x half> @llvm.nvvm.fma.rn.ftz.relu.f16x2(<2 x half> %0, <2 x half> %1, <2 x half> %8)
+ // CHECK-NEXT: %10 = call <2 x half> @llvm.nvvm.fma.rn.oob.v2f16(<2 x half> %0, <2 x half> %1, <2 x half> %9)
+ // CHECK-NEXT: %11 = call <2 x half> @llvm.nvvm.fma.rn.oob.relu.v2f16(<2 x half> %0, <2 x half> %1, <2 x half> %10)
+ // CHECK-NEXT: ret <2 x half> %11
+ // CHECK-NEXT: }
+ %f0 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rn>} : vector<2xf16>
+ %f1 = nvvm.fma %a, %b, %f0 {rnd = #nvvm.fp_rnd_mode<rn>, ftz = true} : vector<2xf16>
+ %f2 = nvvm.fma %a, %b, %f1 {rnd = #nvvm.fp_rnd_mode<rn>, sat = #nvvm.sat_mode<sat>} : vector<2xf16>
+ %f3 = nvvm.fma %a, %b, %f2 {rnd = #nvvm.fp_rnd_mode<rn>, sat = #nvvm.sat_mode<sat>, ftz = true} : vector<2xf16>
+ %f4 = nvvm.fma %a, %b, %f3 {rnd = #nvvm.fp_rnd_mode<rn>, relu = true} : vector<2xf16>
+ %f5 = nvvm.fma %a, %b, %f4 {rnd = #nvvm.fp_rnd_mode<rn>, relu = true, ftz = true} : vector<2xf16>
+ %f6 = nvvm.fma %a, %b, %f5 {rnd = #nvvm.fp_rnd_mode<rn>, oob = true} : vector<2xf16>
+ %f7 = nvvm.fma %a, %b, %f6 {rnd = #nvvm.fp_rnd_mode<rn>, oob = true, relu = true} : vector<2xf16>
+ llvm.return %f7 : vector<2xf16>
+}
+
+llvm.func @fma_bf16(%a: vector<2xbf16>, %b: vector<2xbf16>, %c: vector<2xbf16>) -> vector<2xbf16> {
+ // CHECK-LABEL: define <2 x bfloat> @fma_bf16(<2 x bfloat> %0, <2 x bfloat> %1, <2 x bfloat> %2) {
+ // CHECK-NEXT: %4 = call <2 x bfloat> @llvm.nvvm.fma.rn.bf16x2(<2 x bfloat> %0, <2 x bfloat> %1, <2 x bfloat> %2)
+ // CHECK-NEXT: %5 = call <2 x bfloat> @llvm.nvvm.fma.rn.relu.bf16x2(<2 x bfloat> %0, <2 x bfloat> %1, <2 x bfloat> %4)
+ // CHECK-NEXT: %6 = call <2 x bfloat> @llvm.nvvm.fma.rn.oob.v2bf16(<2 x bfloat> %0, <2 x bfloat> %1, <2 x bfloat> %5)
+ // CHECK-NEXT: %7 = call <2 x bfloat> @llvm.nvvm.fma.rn.oob.relu.v2bf16(<2 x bfloat> %0, <2 x bfloat> %1, <2 x bfloat> %6)
+ // CHECK-NEXT: ret <2 x bfloat> %7
+ // CHECK-NEXT: }
+ %f0 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rn>} : vector<2xbf16>
+ %f1 = nvvm.fma %a, %b, %f0 {rnd = #nvvm.fp_rnd_mode<rn>, relu = true} : vector<2xbf16>
+ %f2 = nvvm.fma %a, %b, %f1 {rnd = #nvvm.fp_rnd_mode<rn>, oob = true} : vector<2xbf16>
+ %f3 = nvvm.fma %a, %b, %f2 {rnd = #nvvm.fp_rnd_mode<rn>, oob = true, relu = true} : vector<2xbf16>
+ llvm.return %f3 : vector<2xbf16>
+}
+
+llvm.func @fma_f32_rn(%a: vector<2xf32>, %b: vector<2xf32>, %c: vector<2xf32>) -> vector<2xf32> {
+ // CHECK-LABEL: define <2 x float> @fma_f32_rn(<2 x float> %0, <2 x float> %1, <2 x float> %2) {
+ // CHECK-NEXT: %4 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %5 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %6 = extractelement <2 x float> %2, i32 0
+ // CHECK-NEXT: %7 = call float @llvm.nvvm.fma.rn.f(float %4, float %5, float %6)
+ // CHECK-NEXT: %8 = insertelement <2 x float> poison, float %7, i32 0
+ // CHECK-NEXT: %9 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %10 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %11 = extractelement <2 x float> %2, i32 1
+ // CHECK-NEXT: %12 = call float @llvm.nvvm.fma.rn.f(float %9, float %10, float %11)
+ // CHECK-NEXT: %13 = insertelement <2 x float> %8, float %12, i32 1
+ // CHECK-NEXT: %14 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %15 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %16 = extractelement <2 x float> %13, i32 0
+ // CHECK-NEXT: %17 = call float @llvm.nvvm.fma.rn.ftz.f(float %14, float %15, float %16)
+ // CHECK-NEXT: %18 = insertelement <2 x float> poison, float %17, i32 0
+ // CHECK-NEXT: %19 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %20 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %21 = extractelement <2 x float> %13, i32 1
+ // CHECK-NEXT: %22 = call float @llvm.nvvm.fma.rn.ftz.f(float %19, float %20, float %21)
+ // CHECK-NEXT: %23 = insertelement <2 x float> %18, float %22, i32 1
+ // CHECK-NEXT: %24 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %25 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %26 = extractelement <2 x float> %23, i32 0
+ // CHECK-NEXT: %27 = call float @llvm.nvvm.fma.rn.sat.f(float %24, float %25, float %26)
+ // CHECK-NEXT: %28 = insertelement <2 x float> poison, float %27, i32 0
+ // CHECK-NEXT: %29 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %30 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %31 = extractelement <2 x float> %23, i32 1
+ // CHECK-NEXT: %32 = call float @llvm.nvvm.fma.rn.sat.f(float %29, float %30, float %31)
+ // CHECK-NEXT: %33 = insertelement <2 x float> %28, float %32, i32 1
+ // CHECK-NEXT: %34 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %35 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %36 = extractelement <2 x float> %33, i32 0
+ // CHECK-NEXT: %37 = call float @llvm.nvvm.fma.rn.ftz.sat.f(float %34, float %35, float %36)
+ // CHECK-NEXT: %38 = insertelement <2 x float> poison, float %37, i32 0
+ // CHECK-NEXT: %39 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %40 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %41 = extractelement <2 x float> %33, i32 1
+ // CHECK-NEXT: %42 = call float @llvm.nvvm.fma.rn.ftz.sat.f(float %39, float %40, float %41)
+ // CHECK-NEXT: %43 = insertelement <2 x float> %38, float %42, i32 1
+ // CHECK-NEXT: ret <2 x float> %43
+ // CHECK-NEXT: }
+ %f0 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rn>} : vector<2xf32>
+ %f1 = nvvm.fma %a, %b, %f0 {rnd = #nvvm.fp_rnd_mode<rn>, ftz = true} : vector<2xf32>
+ %f2 = nvvm.fma %a, %b, %f1 {rnd = #nvvm.fp_rnd_mode<rn>, sat = #nvvm.sat_mode<sat>} : vector<2xf32>
+ %f3 = nvvm.fma %a, %b, %f2 {rnd = #nvvm.fp_rnd_mode<rn>, sat = #nvvm.sat_mode<sat>, ftz = true} : vector<2xf32>
+ llvm.return %f3 : vector<2xf32>
+}
+
+llvm.func @fma_f32_rm(%a: vector<2xf32>, %b: vector<2xf32>, %c: vector<2xf32>) -> vector<2xf32> {
+ // CHECK-LABEL: define <2 x float> @fma_f32_rm(<2 x float> %0, <2 x float> %1, <2 x float> %2) {
+ // CHECK-NEXT: %4 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %5 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %6 = extractelement <2 x float> %2, i32 0
+ // CHECK-NEXT: %7 = call float @llvm.nvvm.fma.rm.f(float %4, float %5, float %6)
+ // CHECK-NEXT: %8 = insertelement <2 x float> poison, float %7, i32 0
+ // CHECK-NEXT: %9 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %10 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %11 = extractelement <2 x float> %2, i32 1
+ // CHECK-NEXT: %12 = call float @llvm.nvvm.fma.rm.f(float %9, float %10, float %11)
+ // CHECK-NEXT: %13 = insertelement <2 x float> %8, float %12, i32 1
+ // CHECK-NEXT: %14 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %15 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %16 = extractelement <2 x float> %13, i32 0
+ // CHECK-NEXT: %17 = call float @llvm.nvvm.fma.rm.ftz.f(float %14, float %15, float %16)
+ // CHECK-NEXT: %18 = insertelement <2 x float> poison, float %17, i32 0
+ // CHECK-NEXT: %19 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %20 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %21 = extractelement <2 x float> %13, i32 1
+ // CHECK-NEXT: %22 = call float @llvm.nvvm.fma.rm.ftz.f(float %19, float %20, float %21)
+ // CHECK-NEXT: %23 = insertelement <2 x float> %18, float %22, i32 1
+ // CHECK-NEXT: %24 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %25 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %26 = extractelement <2 x float> %23, i32 0
+ // CHECK-NEXT: %27 = call float @llvm.nvvm.fma.rm.sat.f(float %24, float %25, float %26)
+ // CHECK-NEXT: %28 = insertelement <2 x float> poison, float %27, i32 0
+ // CHECK-NEXT: %29 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %30 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %31 = extractelement <2 x float> %23, i32 1
+ // CHECK-NEXT: %32 = call float @llvm.nvvm.fma.rm.sat.f(float %29, float %30, float %31)
+ // CHECK-NEXT: %33 = insertelement <2 x float> %28, float %32, i32 1
+ // CHECK-NEXT: %34 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %35 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %36 = extractelement <2 x float> %33, i32 0
+ // CHECK-NEXT: %37 = call float @llvm.nvvm.fma.rm.ftz.sat.f(float %34, float %35, float %36)
+ // CHECK-NEXT: %38 = insertelement <2 x float> poison, float %37, i32 0
+ // CHECK-NEXT: %39 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %40 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %41 = extractelement <2 x float> %33, i32 1
+ // CHECK-NEXT: %42 = call float @llvm.nvvm.fma.rm.ftz.sat.f(float %39, float %40, float %41)
+ // CHECK-NEXT: %43 = insertelement <2 x float> %38, float %42, i32 1
+ // CHECK-NEXT: ret <2 x float> %43
+ // CHECK-NEXT: }
+ %f0 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rm>} : vector<2xf32>
+ %f1 = nvvm.fma %a, %b, %f0 {rnd = #nvvm.fp_rnd_mode<rm>, ftz = true} : vector<2xf32>
+ %f2 = nvvm.fma %a, %b, %f1 {rnd = #nvvm.fp_rnd_mode<rm>, sat = #nvvm.sat_mode<sat>} : vector<2xf32>
+ %f3 = nvvm.fma %a, %b, %f2 {rnd = #nvvm.fp_rnd_mode<rm>, sat = #nvvm.sat_mode<sat>, ftz = true} : vector<2xf32>
+ llvm.return %f3 : vector<2xf32>
+}
+
+llvm.func @fma_f32_rp(%a: vector<2xf32>, %b: vector<2xf32>, %c: vector<2xf32>) -> vector<2xf32> {
+ // CHECK-LABEL: define <2 x float> @fma_f32_rp(<2 x float> %0, <2 x float> %1, <2 x float> %2) {
+ // CHECK-NEXT: %4 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %5 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %6 = extractelement <2 x float> %2, i32 0
+ // CHECK-NEXT: %7 = call float @llvm.nvvm.fma.rp.f(float %4, float %5, float %6)
+ // CHECK-NEXT: %8 = insertelement <2 x float> poison, float %7, i32 0
+ // CHECK-NEXT: %9 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %10 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %11 = extractelement <2 x float> %2, i32 1
+ // CHECK-NEXT: %12 = call float @llvm.nvvm.fma.rp.f(float %9, float %10, float %11)
+ // CHECK-NEXT: %13 = insertelement <2 x float> %8, float %12, i32 1
+ // CHECK-NEXT: %14 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %15 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %16 = extractelement <2 x float> %13, i32 0
+ // CHECK-NEXT: %17 = call float @llvm.nvvm.fma.rp.ftz.f(float %14, float %15, float %16)
+ // CHECK-NEXT: %18 = insertelement <2 x float> poison, float %17, i32 0
+ // CHECK-NEXT: %19 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %20 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %21 = extractelement <2 x float> %13, i32 1
+ // CHECK-NEXT: %22 = call float @llvm.nvvm.fma.rp.ftz.f(float %19, float %20, float %21)
+ // CHECK-NEXT: %23 = insertelement <2 x float> %18, float %22, i32 1
+ // CHECK-NEXT: %24 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %25 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %26 = extractelement <2 x float> %23, i32 0
+ // CHECK-NEXT: %27 = call float @llvm.nvvm.fma.rp.sat.f(float %24, float %25, float %26)
+ // CHECK-NEXT: %28 = insertelement <2 x float> poison, float %27, i32 0
+ // CHECK-NEXT: %29 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %30 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %31 = extractelement <2 x float> %23, i32 1
+ // CHECK-NEXT: %32 = call float @llvm.nvvm.fma.rp.sat.f(float %29, float %30, float %31)
+ // CHECK-NEXT: %33 = insertelement <2 x float> %28, float %32, i32 1
+ // CHECK-NEXT: %34 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %35 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %36 = extractelement <2 x float> %33, i32 0
+ // CHECK-NEXT: %37 = call float @llvm.nvvm.fma.rp.ftz.sat.f(float %34, float %35, float %36)
+ // CHECK-NEXT: %38 = insertelement <2 x float> poison, float %37, i32 0
+ // CHECK-NEXT: %39 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %40 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %41 = extractelement <2 x float> %33, i32 1
+ // CHECK-NEXT: %42 = call float @llvm.nvvm.fma.rp.ftz.sat.f(float %39, float %40, float %41)
+ // CHECK-NEXT: %43 = insertelement <2 x float> %38, float %42, i32 1
+ // CHECK-NEXT: ret <2 x float> %43
+ // CHECK-NEXT: }
+ %f0 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rp>} : vector<2xf32>
+ %f1 = nvvm.fma %a, %b, %f0 {rnd = #nvvm.fp_rnd_mode<rp>, ftz = true} : vector<2xf32>
+ %f2 = nvvm.fma %a, %b, %f1 {rnd = #nvvm.fp_rnd_mode<rp>, sat = #nvvm.sat_mode<sat>} : vector<2xf32>
+ %f3 = nvvm.fma %a, %b, %f2 {rnd = #nvvm.fp_rnd_mode<rp>, sat = #nvvm.sat_mode<sat>, ftz = true} : vector<2xf32>
+ llvm.return %f3 : vector<2xf32>
+}
+
+llvm.func @fma_f32_rz(%a: vector<2xf32>, %b: vector<2xf32>, %c: vector<2xf32>) -> vector<2xf32> {
+ // CHECK-LABEL: define <2 x float> @fma_f32_rz(<2 x float> %0, <2 x float> %1, <2 x float> %2) {
+ // CHECK-NEXT: %4 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %5 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %6 = extractelement <2 x float> %2, i32 0
+ // CHECK-NEXT: %7 = call float @llvm.nvvm.fma.rz.f(float %4, float %5, float %6)
+ // CHECK-NEXT: %8 = insertelement <2 x float> poison, float %7, i32 0
+ // CHECK-NEXT: %9 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %10 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %11 = extractelement <2 x float> %2, i32 1
+ // CHECK-NEXT: %12 = call float @llvm.nvvm.fma.rz.f(float %9, float %10, float %11)
+ // CHECK-NEXT: %13 = insertelement <2 x float> %8, float %12, i32 1
+ // CHECK-NEXT: %14 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %15 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %16 = extractelement <2 x float> %13, i32 0
+ // CHECK-NEXT: %17 = call float @llvm.nvvm.fma.rz.ftz.f(float %14, float %15, float %16)
+ // CHECK-NEXT: %18 = insertelement <2 x float> poison, float %17, i32 0
+ // CHECK-NEXT: %19 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %20 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %21 = extractelement <2 x float> %13, i32 1
+ // CHECK-NEXT: %22 = call float @llvm.nvvm.fma.rz.ftz.f(float %19, float %20, float %21)
+ // CHECK-NEXT: %23 = insertelement <2 x float> %18, float %22, i32 1
+ // CHECK-NEXT: %24 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %25 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %26 = extractelement <2 x float> %23, i32 0
+ // CHECK-NEXT: %27 = call float @llvm.nvvm.fma.rz.sat.f(float %24, float %25, float %26)
+ // CHECK-NEXT: %28 = insertelement <2 x float> poison, float %27, i32 0
+ // CHECK-NEXT: %29 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %30 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %31 = extractelement <2 x float> %23, i32 1
+ // CHECK-NEXT: %32 = call float @llvm.nvvm.fma.rz.sat.f(float %29, float %30, float %31)
+ // CHECK-NEXT: %33 = insertelement <2 x float> %28, float %32, i32 1
+ // CHECK-NEXT: %34 = extractelement <2 x float> %0, i32 0
+ // CHECK-NEXT: %35 = extractelement <2 x float> %1, i32 0
+ // CHECK-NEXT: %36 = extractelement <2 x float> %33, i32 0
+ // CHECK-NEXT: %37 = call float @llvm.nvvm.fma.rz.ftz.sat.f(float %34, float %35, float %36)
+ // CHECK-NEXT: %38 = insertelement <2 x float> poison, float %37, i32 0
+ // CHECK-NEXT: %39 = extractelement <2 x float> %0, i32 1
+ // CHECK-NEXT: %40 = extractelement <2 x float> %1, i32 1
+ // CHECK-NEXT: %41 = extractelement <2 x float> %33, i32 1
+ // CHECK-NEXT: %42 = call float @llvm.nvvm.fma.rz.ftz.sat.f(float %39, float %40, float %41)
+ // CHECK-NEXT: %43 = insertelement <2 x float> %38, float %42, i32 1
+ // CHECK-NEXT: ret <2 x float> %43
+ // CHECK-NEXT: }
+ %f0 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rz>} : vector<2xf32>
+ %f1 = nvvm.fma %a, %b, %f0 {rnd = #nvvm.fp_rnd_mode<rz>, ftz = true} : vector<2xf32>
+ %f2 = nvvm.fma %a, %b, %f1 {rnd = #nvvm.fp_rnd_mode<rz>, sat = #nvvm.sat_mode<sat>} : vector<2xf32>
+ %f3 = nvvm.fma %a, %b, %f2 {rnd = #nvvm.fp_rnd_mode<rz>, sat = #nvvm.sat_mode<sat>, ftz = true} : vector<2xf32>
+ llvm.return %f3 : vector<2xf32>
+}
+
+llvm.func @fma_f64(%a: vector<2xf64>, %b: vector<2xf64>, %c: vector<2xf64>) -> vector<2xf64> {
+ // CHECK-LABEL: define <2 x double> @fma_f64(<2 x double> %0, <2 x double> %1, <2 x double> %2) {
+ // CHECK-NEXT: %4 = extractelement <2 x double> %0, i32 0
+ // CHECK-NEXT: %5 = extractelement <2 x double> %1, i32 0
+ // CHECK-NEXT: %6 = extractelement <2 x double> %2, i32 0
+ // CHECK-NEXT: %7 = call double @llvm.nvvm.fma.rn.d(double %4, double %5, double %6)
+ // CHECK-NEXT: %8 = insertelement <2 x double> poison, double %7, i32 0
+ // CHECK-NEXT: %9 = extractelement <2 x double> %0, i32 1
+ // CHECK-NEXT: %10 = extractelement <2 x double> %1, i32 1
+ // CHECK-NEXT: %11 = extractelement <2 x double> %2, i32 1
+ // CHECK-NEXT: %12 = call double @llvm.nvvm.fma.rn.d(double %9, double %10, double %11)
+ // CHECK-NEXT: %13 = insertelement <2 x double> %8, double %12, i32 1
+ // CHECK-NEXT: %14 = extractelement <2 x double> %0, i32 0
+ // CHECK-NEXT: %15 = extractelement <2 x double> %1, i32 0
+ // CHECK-NEXT: %16 = extractelement <2 x double> %13, i32 0
+ // CHECK-NEXT: %17 = call double @llvm.nvvm.fma.rm.d(double %14, double %15, double %16)
+ // CHECK-NEXT: %18 = insertelement <2 x double> poison, double %17, i32 0
+ // CHECK-NEXT: %19 = extractelement <2 x double> %0, i32 1
+ // CHECK-NEXT: %20 = extractelement <2 x double> %1, i32 1
+ // CHECK-NEXT: %21 = extractelement <2 x double> %13, i32 1
+ // CHECK-NEXT: %22 = call double @llvm.nvvm.fma.rm.d(double %19, double %20, double %21)
+ // CHECK-NEXT: %23 = insertelement <2 x double> %18, double %22, i32 1
+ // CHECK-NEXT: %24 = extractelement <2 x double> %0, i32 0
+ // CHECK-NEXT: %25 = extractelement <2 x double> %1, i32 0
+ // CHECK-NEXT: %26 = extractelement <2 x double> %23, i32 0
+ // CHECK-NEXT: %27 = call double @llvm.nvvm.fma.rp.d(double %24, double %25, double %26)
+ // CHECK-NEXT: %28 = insertelement <2 x double> poison, double %27, i32 0
+ // CHECK-NEXT: %29 = extractelement <2 x double> %0, i32 1
+ // CHECK-NEXT: %30 = extractelement <2 x double> %1, i32 1
+ // CHECK-NEXT: %31 = extractelement <2 x double> %23, i32 1
+ // CHECK-NEXT: %32 = call double @llvm.nvvm.fma.rp.d(double %29, double %30, double %31)
+ // CHECK-NEXT: %33 = insertelement <2 x double> %28, double %32, i32 1
+ // CHECK-NEXT: %34 = extractelement <2 x double> %0, i32 0
+ // CHECK-NEXT: %35 = extractelement <2 x double> %1, i32 0
+ // CHECK-NEXT: %36 = extractelement <2 x double> %33, i32 0
+ // CHECK-NEXT: %37 = call double @llvm.nvvm.fma.rz.d(double %34, double %35, double %36)
+ // CHECK-NEXT: %38 = insertelement <2 x double> poison, double %37, i32 0
+ // CHECK-NEXT: %39 = extractelement <2 x double> %0, i32 1
+ // CHECK-NEXT: %40 = extractelement <2 x double> %1, i32 1
+ // CHECK-NEXT: %41 = extractelement <2 x double> %33, i32 1
+ // CHECK-NEXT: %42 = call double @llvm.nvvm.fma.rz.d(double %39, double %40, double %41)
+ // CHECK-NEXT: %43 = insertelement <2 x double> %38, double %42, i32 1
+ // CHECK-NEXT: ret <2 x double> %43
+ // CHECK-NEXT: }
+ %f0 = nvvm.fma %a, %b, %c {rnd = #nvvm.fp_rnd_mode<rn>} : vector<2xf64>
+ %f1 = nvvm.fma %a, %b, %f0 {rnd = #nvvm.fp_rnd_mode<rm>} : vector<2xf64>
+ %f2 = nvvm.fma %a, %b, %f1 {rnd = #nvvm.fp_rnd_mode<rp>} : vector<2xf64>
+ %f3 = nvvm.fma %a, %b, %f2 {rnd = #nvvm.fp_rnd_mode<rz>} : vector<2xf64>
+ llvm.return %f3 : vector<2xf64>
+}
More information about the Mlir-commits
mailing list