[Mlir-commits] [clang] [llvm] [mlir] [clang][NVPTX][MLIR][NVVM] Add overloaded fma intrinsics (PR #232672)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Fri Oct 16 03:24:17 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-llvm-transforms
@llvm/pr-subscribers-backend-nvptx
Author: Srinivasa Ravi (Wolfram70)
<details>
<summary>Changes</summary>
This change adds the following overloaded `fma` intrinsics with
NVPTX codegen:
- `llvm.nvvm.fma`
- `llvm.nvvm.fma.ftz`
- `llvm.nvvm.fma.sat`
- `llvm.nvvm.fma.ftz.sat`
- `llvm.nvvm.fma.relu`
- `llvm.nvvm.fma.ftz.relu`
The rounding mode is passed in as an `i32` immediate operand.
Auto-upgrades the older non-overloaded intrinsics (including the
legacy variants that used integer storage types for the bf16
operands) to the new ones, and updates clang builtins, CIR codegen,
and MLIR NVVM ops to lower to the new intrinsics. The
`llvm.nvvm.fma.rn.oob{.relu}` intrinsics are left unchanged.
In the interest of completion, this also:
- Adds support for the `f32x2` type (`fma.rnd{.ftz}.f32x2`,
`sm_100+`).
- Adds tests for constant folding of these intrinsics with the newly
supported scalar types.
PTX Spec References:
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-fma
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-fma
Assisted-by: Claude Opus 5.5
---
Patch is 202.22 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/232672.diff
30 Files Affected:
- (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp (+81-30)
- (modified) clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp (+64-20)
- (modified) clang/test/CIR/CodeGenCUDA/builtins-nvvm-half.cu (+24-24)
- (modified) clang/test/CIR/CodeGenCUDA/builtins-nvvm-math.cu (+32)
- (modified) clang/test/CodeGen/builtins-nvptx-native-half-type-native.c (+11-11)
- (modified) clang/test/CodeGen/builtins-nvptx-native-half-type.c (+12-12)
- (modified) clang/test/CodeGen/builtins-nvptx.c (+12-12)
- (modified) llvm/docs/NVPTXUsage.md (+51-18)
- (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+17-37)
- (modified) llvm/include/llvm/IR/NVVMIntrinsicUtils.h (+40-45)
- (modified) llvm/lib/Analysis/ConstantFolding.cpp (+47-46)
- (modified) llvm/lib/IR/AutoUpgrade.cpp (+60-8)
- (modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp (+16-2)
- (modified) llvm/lib/Target/NVPTX/NVPTXIntrinsics.td (+65-56)
- (modified) llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp (+21-18)
- (modified) llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll (+57)
- (modified) llvm/test/CodeGen/NVPTX/fp-arith-sat.ll (+32-32)
- (added) llvm/test/CodeGen/NVPTX/fp-fma-f32x2.ll (+133)
- (added) llvm/test/CodeGen/NVPTX/fp-fma-invalid.ll (+65)
- (modified) llvm/test/CodeGen/NVPTX/math-intrins-sm53-ptx42.ll (+16-16)
- (modified) llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70-instcombine.ll (+16-16)
- (modified) llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70.ll (+16-16)
- (modified) llvm/test/CodeGen/NVPTX/mixed-precision-fp.ll (+48-48)
- (modified) llvm/test/Transforms/InstCombine/NVPTX/nvvm-intrins.ll (+8-8)
- (modified) llvm/test/Transforms/InstSimplify/const-fold-nvvm-fma.ll (+380-108)
- (added) llvm/test/Verifier/NVPTX/fma.ll (+16)
- (modified) llvm/test/Verifier/invalid_intrinsic_signatures.ll (+2-2)
- (modified) mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp (+26-92)
- (modified) mlir/test/Target/LLVMIR/nvvm/fma/fma.mlir (+28-28)
- (modified) mlir/test/Target/LLVMIR/nvvm/fma/fma_vector.mlir (+108-180)
``````````diff
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
index 146dad6b0ec413..fc984df8e7d67b 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
@@ -85,6 +85,24 @@ static mlir::Value emitNVVMFPArith(CIRGenFunction &cgf, const CallExpr *expr,
.getResult();
}
+/// Emit a CIR LLVMIntrinsicCallOp for an NVVM fma intrinsic, which takes the
+/// rounding mode as a trailing operand.
+static mlir::Value emitNVVMFMA(CIRGenFunction &cgf, const CallExpr *expr,
+ llvm::StringRef intrinsicName,
+ llvm::APFloat::roundingMode rm) {
+ auto &builder = cgf.getBuilder();
+ mlir::Location loc = cgf.getLoc(expr->getExprLoc());
+ mlir::Value a = cgf.emitScalarExpr(expr->getArg(0));
+ mlir::Value b = cgf.emitScalarExpr(expr->getArg(1));
+ mlir::Value c = cgf.emitScalarExpr(expr->getArg(2));
+ mlir::Value rnd =
+ builder.getConstInt(loc, builder.getSInt32Ty(), static_cast<int>(rm));
+ return cir::LLVMIntrinsicCallOp::create(builder, loc,
+ builder.getStringAttr(intrinsicName),
+ a.getType(), {a, b, c, rnd})
+ .getResult();
+}
+
static mlir::Value emitBar0Reduction(CIRGenFunction &cgf, const CallExpr *expr,
llvm::StringRef intrinsicName,
bool returnsPred) {
@@ -552,43 +570,76 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) {
getContext().BuiltinInfo.getName(builtinId));
return mlir::Value{};
case NVPTX::BI__nvvm_fma_rn_f16:
- return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.f16")
- .getValue();
case NVPTX::BI__nvvm_fma_rn_f16x2:
- return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.f16x2")
- .getValue();
+ case NVPTX::BI__nvvm_fma_rn_bf16:
+ case NVPTX::BI__nvvm_fma_rn_bf16x2:
+ case NVPTX::BI__nvvm_fma_rn_f:
+ case NVPTX::BI__nvvm_fma_rn_d:
+ return emitNVVMFMA(*this, expr, "nvvm.fma",
+ llvm::APFloat::rmNearestTiesToEven);
+ case NVPTX::BI__nvvm_fma_rz_f:
+ case NVPTX::BI__nvvm_fma_rz_d:
+ return emitNVVMFMA(*this, expr, "nvvm.fma", llvm::APFloat::rmTowardZero);
+ case NVPTX::BI__nvvm_fma_rm_f:
+ case NVPTX::BI__nvvm_fma_rm_d:
+ return emitNVVMFMA(*this, expr, "nvvm.fma",
+ llvm::APFloat::rmTowardNegative);
+ case NVPTX::BI__nvvm_fma_rp_f:
+ case NVPTX::BI__nvvm_fma_rp_d:
+ return emitNVVMFMA(*this, expr, "nvvm.fma",
+ llvm::APFloat::rmTowardPositive);
case NVPTX::BI__nvvm_fma_rn_ftz_f16:
- return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.ftz.f16")
- .getValue();
case NVPTX::BI__nvvm_fma_rn_ftz_f16x2:
- return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.ftz.f16x2")
- .getValue();
- case NVPTX::BI__nvvm_fma_rn_ftz_relu_f16:
- return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.ftz.relu.f16")
- .getValue();
- case NVPTX::BI__nvvm_fma_rn_ftz_relu_f16x2:
- return emitBuiltinWithOneOverloadedType<3>(expr,
- "nvvm.fma.rn.ftz.relu.f16x2")
- .getValue();
+ case NVPTX::BI__nvvm_fma_rn_ftz_f:
+ return emitNVVMFMA(*this, expr, "nvvm.fma.ftz",
+ llvm::APFloat::rmNearestTiesToEven);
+ case NVPTX::BI__nvvm_fma_rz_ftz_f:
+ return emitNVVMFMA(*this, expr, "nvvm.fma.ftz",
+ llvm::APFloat::rmTowardZero);
+ case NVPTX::BI__nvvm_fma_rm_ftz_f:
+ return emitNVVMFMA(*this, expr, "nvvm.fma.ftz",
+ llvm::APFloat::rmTowardNegative);
+ case NVPTX::BI__nvvm_fma_rp_ftz_f:
+ return emitNVVMFMA(*this, expr, "nvvm.fma.ftz",
+ llvm::APFloat::rmTowardPositive);
+ case NVPTX::BI__nvvm_fma_rn_sat_f16:
+ case NVPTX::BI__nvvm_fma_rn_sat_f16x2:
+ case NVPTX::BI__nvvm_fma_rn_sat_f:
+ return emitNVVMFMA(*this, expr, "nvvm.fma.sat",
+ llvm::APFloat::rmNearestTiesToEven);
+ case NVPTX::BI__nvvm_fma_rz_sat_f:
+ return emitNVVMFMA(*this, expr, "nvvm.fma.sat",
+ llvm::APFloat::rmTowardZero);
+ case NVPTX::BI__nvvm_fma_rm_sat_f:
+ return emitNVVMFMA(*this, expr, "nvvm.fma.sat",
+ llvm::APFloat::rmTowardNegative);
+ case NVPTX::BI__nvvm_fma_rp_sat_f:
+ return emitNVVMFMA(*this, expr, "nvvm.fma.sat",
+ llvm::APFloat::rmTowardPositive);
case NVPTX::BI__nvvm_fma_rn_ftz_sat_f16:
- return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.ftz.sat.f16")
- .getValue();
case NVPTX::BI__nvvm_fma_rn_ftz_sat_f16x2:
- return emitBuiltinWithOneOverloadedType<3>(expr,
- "nvvm.fma.rn.ftz.sat.f16x2")
- .getValue();
+ case NVPTX::BI__nvvm_fma_rn_ftz_sat_f:
+ return emitNVVMFMA(*this, expr, "nvvm.fma.ftz.sat",
+ llvm::APFloat::rmNearestTiesToEven);
+ case NVPTX::BI__nvvm_fma_rz_ftz_sat_f:
+ return emitNVVMFMA(*this, expr, "nvvm.fma.ftz.sat",
+ llvm::APFloat::rmTowardZero);
+ case NVPTX::BI__nvvm_fma_rm_ftz_sat_f:
+ return emitNVVMFMA(*this, expr, "nvvm.fma.ftz.sat",
+ llvm::APFloat::rmTowardNegative);
+ case NVPTX::BI__nvvm_fma_rp_ftz_sat_f:
+ return emitNVVMFMA(*this, expr, "nvvm.fma.ftz.sat",
+ llvm::APFloat::rmTowardPositive);
case NVPTX::BI__nvvm_fma_rn_relu_f16:
- return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.relu.f16")
- .getValue();
case NVPTX::BI__nvvm_fma_rn_relu_f16x2:
- return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.relu.f16x2")
- .getValue();
- case NVPTX::BI__nvvm_fma_rn_sat_f16:
- return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.sat.f16")
- .getValue();
- case NVPTX::BI__nvvm_fma_rn_sat_f16x2:
- return emitBuiltinWithOneOverloadedType<3>(expr, "nvvm.fma.rn.sat.f16x2")
- .getValue();
+ case NVPTX::BI__nvvm_fma_rn_relu_bf16:
+ case NVPTX::BI__nvvm_fma_rn_relu_bf16x2:
+ return emitNVVMFMA(*this, expr, "nvvm.fma.relu",
+ llvm::APFloat::rmNearestTiesToEven);
+ case NVPTX::BI__nvvm_fma_rn_ftz_relu_f16:
+ case NVPTX::BI__nvvm_fma_rn_ftz_relu_f16x2:
+ return emitNVVMFMA(*this, expr, "nvvm.fma.ftz.relu",
+ llvm::APFloat::rmNearestTiesToEven);
case NVPTX::BI__nvvm_fma_rn_oob_f16:
case NVPTX::BI__nvvm_fma_rn_oob_f16x2:
case NVPTX::BI__nvvm_fma_rn_oob_bf16:
diff --git a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
index a513fcc12953ba..691d181d89bb8a 100644
--- a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
@@ -1180,33 +1180,77 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
#undef PZO_CVT
case NVPTX::BI__nvvm_fma_rn_f16:
- return MakeHalfType(Intrinsic::nvvm_fma_rn_f16, BuiltinID, E, *this);
case NVPTX::BI__nvvm_fma_rn_f16x2:
- return MakeHalfType(Intrinsic::nvvm_fma_rn_f16x2, BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_fma_rn_bf16:
+ case NVPTX::BI__nvvm_fma_rn_bf16x2:
+ case NVPTX::BI__nvvm_fma_rn_f:
+ case NVPTX::BI__nvvm_fma_rn_d:
+ return MakeFPArith(Intrinsic::nvvm_fma, APFloat::rmNearestTiesToEven,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_fma_rz_f:
+ case NVPTX::BI__nvvm_fma_rz_d:
+ return MakeFPArith(Intrinsic::nvvm_fma, APFloat::rmTowardZero, BuiltinID, E,
+ *this);
+ case NVPTX::BI__nvvm_fma_rm_f:
+ case NVPTX::BI__nvvm_fma_rm_d:
+ return MakeFPArith(Intrinsic::nvvm_fma, APFloat::rmTowardNegative,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_fma_rp_f:
+ case NVPTX::BI__nvvm_fma_rp_d:
+ return MakeFPArith(Intrinsic::nvvm_fma, APFloat::rmTowardPositive,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_fma_rn_ftz_f16:
- return MakeHalfType(Intrinsic::nvvm_fma_rn_ftz_f16, BuiltinID, E, *this);
case NVPTX::BI__nvvm_fma_rn_ftz_f16x2:
- return MakeHalfType(Intrinsic::nvvm_fma_rn_ftz_f16x2, BuiltinID, E, *this);
- case NVPTX::BI__nvvm_fma_rn_ftz_relu_f16:
- return MakeHalfType(Intrinsic::nvvm_fma_rn_ftz_relu_f16, BuiltinID, E,
- *this);
- case NVPTX::BI__nvvm_fma_rn_ftz_relu_f16x2:
- return MakeHalfType(Intrinsic::nvvm_fma_rn_ftz_relu_f16x2, BuiltinID, E,
- *this);
+ case NVPTX::BI__nvvm_fma_rn_ftz_f:
+ return MakeFPArith(Intrinsic::nvvm_fma_ftz, APFloat::rmNearestTiesToEven,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_fma_rz_ftz_f:
+ return MakeFPArith(Intrinsic::nvvm_fma_ftz, APFloat::rmTowardZero,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_fma_rm_ftz_f:
+ return MakeFPArith(Intrinsic::nvvm_fma_ftz, APFloat::rmTowardNegative,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_fma_rp_ftz_f:
+ return MakeFPArith(Intrinsic::nvvm_fma_ftz, APFloat::rmTowardPositive,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_fma_rn_sat_f16:
+ case NVPTX::BI__nvvm_fma_rn_sat_f16x2:
+ case NVPTX::BI__nvvm_fma_rn_sat_f:
+ return MakeFPArith(Intrinsic::nvvm_fma_sat, APFloat::rmNearestTiesToEven,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_fma_rz_sat_f:
+ return MakeFPArith(Intrinsic::nvvm_fma_sat, APFloat::rmTowardZero,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_fma_rm_sat_f:
+ return MakeFPArith(Intrinsic::nvvm_fma_sat, APFloat::rmTowardNegative,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_fma_rp_sat_f:
+ return MakeFPArith(Intrinsic::nvvm_fma_sat, APFloat::rmTowardPositive,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_fma_rn_ftz_sat_f16:
- return MakeHalfType(Intrinsic::nvvm_fma_rn_ftz_sat_f16, BuiltinID, E,
- *this);
case NVPTX::BI__nvvm_fma_rn_ftz_sat_f16x2:
- return MakeHalfType(Intrinsic::nvvm_fma_rn_ftz_sat_f16x2, BuiltinID, E,
- *this);
+ case NVPTX::BI__nvvm_fma_rn_ftz_sat_f:
+ return MakeFPArith(Intrinsic::nvvm_fma_ftz_sat,
+ APFloat::rmNearestTiesToEven, BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_fma_rz_ftz_sat_f:
+ return MakeFPArith(Intrinsic::nvvm_fma_ftz_sat, APFloat::rmTowardZero,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_fma_rm_ftz_sat_f:
+ return MakeFPArith(Intrinsic::nvvm_fma_ftz_sat, APFloat::rmTowardNegative,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_fma_rp_ftz_sat_f:
+ return MakeFPArith(Intrinsic::nvvm_fma_ftz_sat, APFloat::rmTowardPositive,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_fma_rn_relu_f16:
- return MakeHalfType(Intrinsic::nvvm_fma_rn_relu_f16, BuiltinID, E, *this);
case NVPTX::BI__nvvm_fma_rn_relu_f16x2:
- return MakeHalfType(Intrinsic::nvvm_fma_rn_relu_f16x2, BuiltinID, E, *this);
- case NVPTX::BI__nvvm_fma_rn_sat_f16:
- return MakeHalfType(Intrinsic::nvvm_fma_rn_sat_f16, BuiltinID, E, *this);
- case NVPTX::BI__nvvm_fma_rn_sat_f16x2:
- return MakeHalfType(Intrinsic::nvvm_fma_rn_sat_f16x2, BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_fma_rn_relu_bf16:
+ case NVPTX::BI__nvvm_fma_rn_relu_bf16x2:
+ return MakeFPArith(Intrinsic::nvvm_fma_relu, APFloat::rmNearestTiesToEven,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_fma_rn_ftz_relu_f16:
+ case NVPTX::BI__nvvm_fma_rn_ftz_relu_f16x2:
+ return MakeFPArith(Intrinsic::nvvm_fma_ftz_relu,
+ APFloat::rmNearestTiesToEven, BuiltinID, E, *this);
case NVPTX::BI__nvvm_fma_rn_oob_f16:
return MakeFMAOOB(Intrinsic::nvvm_fma_rn_oob, Builder.getHalfTy(), E,
*this);
diff --git a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-half.cu b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-half.cu
index d8dd88616d830c..98d902f8aa9c31 100644
--- a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-half.cu
+++ b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-half.cu
@@ -39,97 +39,97 @@ extern "C" __device__ f16x2 test_ex2_approx_f16x2(f16x2 a0) {
}
// CIR-LABEL: @test_fma_rn_f16
-// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.f16" {{.*}} : (!cir.f16, !cir.f16, !cir.f16) -> !cir.f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma" {{.*}} : (!cir.f16, !cir.f16, !cir.f16, !s32i) -> !cir.f16
// LLVM-LABEL: @test_fma_rn_f16(
-// LLVM: call{{.*}} half @llvm.nvvm.fma.rn.f16(half %{{.*}}, half %{{.*}}, half %{{.*}})
+// LLVM: call{{.*}} half @llvm.nvvm.fma.f16(half %{{.*}}, half %{{.*}}, half %{{.*}}, /* rnd=rn */ i32 1)
extern "C" __device__ _Float16 test_fma_rn_f16(_Float16 a0, _Float16 a1, _Float16 a2) {
return __nvvm_fma_rn_f16(a0, a1, a2);
}
// CIR-LABEL: @test_fma_rn_f16x2
-// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.f16x2" {{.*}} : (!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x !cir.f16>
+// CIR: cir.call_llvm_intrinsic "nvvm.fma" {{.*}} : (!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !s32i) -> !cir.vector<2 x !cir.f16>
// LLVM-LABEL: @test_fma_rn_f16x2(
-// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.rn.f16x2(<2 x half> %{{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}})
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.v2f16(<2 x half> %{{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}}, /* rnd=rn */ i32 1)
extern "C" __device__ f16x2 test_fma_rn_f16x2(f16x2 a0, f16x2 a1, f16x2 a2) {
return __nvvm_fma_rn_f16x2(a0, a1, a2);
}
// CIR-LABEL: @test_fma_rn_ftz_f16
-// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.ftz.f16" {{.*}} : (!cir.f16, !cir.f16, !cir.f16) -> !cir.f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.ftz" {{.*}} : (!cir.f16, !cir.f16, !cir.f16, !s32i) -> !cir.f16
// LLVM-LABEL: @test_fma_rn_ftz_f16(
-// LLVM: call{{.*}} half @llvm.nvvm.fma.rn.ftz.f16(half %{{.*}}, half %{{.*}}, half %{{.*}})
+// LLVM: call{{.*}} half @llvm.nvvm.fma.ftz.f16(half %{{.*}}, half %{{.*}}, half %{{.*}}, /* rnd=rn */ i32 1)
extern "C" __device__ _Float16 test_fma_rn_ftz_f16(_Float16 a0, _Float16 a1, _Float16 a2) {
return __nvvm_fma_rn_ftz_f16(a0, a1, a2);
}
// CIR-LABEL: @test_fma_rn_ftz_f16x2
-// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.ftz.f16x2" {{.*}} : (!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x !cir.f16>
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.ftz" {{.*}} : (!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !s32i) -> !cir.vector<2 x !cir.f16>
// LLVM-LABEL: @test_fma_rn_ftz_f16x2(
-// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.rn.ftz.f16x2(<2 x half> %{{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}})
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.ftz.v2f16(<2 x half> %{{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}}, /* rnd=rn */ i32 1)
extern "C" __device__ f16x2 test_fma_rn_ftz_f16x2(f16x2 a0, f16x2 a1, f16x2 a2) {
return __nvvm_fma_rn_ftz_f16x2(a0, a1, a2);
}
// CIR-LABEL: @test_fma_rn_ftz_relu_f16
-// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.ftz.relu.f16" {{.*}} : (!cir.f16, !cir.f16, !cir.f16) -> !cir.f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.ftz.relu" {{.*}} : (!cir.f16, !cir.f16, !cir.f16, !s32i) -> !cir.f16
// LLVM-LABEL: @test_fma_rn_ftz_relu_f16(
-// LLVM: call{{.*}} half @llvm.nvvm.fma.rn.ftz.relu.f16(half %{{.*}}, half %{{.*}}, half %{{.*}})
+// LLVM: call{{.*}} half @llvm.nvvm.fma.ftz.relu.f16(half %{{.*}}, half %{{.*}}, half %{{.*}}, /* rnd=rn */ i32 1)
extern "C" __device__ _Float16 test_fma_rn_ftz_relu_f16(_Float16 a0, _Float16 a1, _Float16 a2) {
return __nvvm_fma_rn_ftz_relu_f16(a0, a1, a2);
}
// CIR-LABEL: @test_fma_rn_ftz_relu_f16x2
-// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.ftz.relu.f16x2" {{.*}} : (!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x !cir.f16>
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.ftz.relu" {{.*}} : (!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !s32i) -> !cir.vector<2 x !cir.f16>
// LLVM-LABEL: @test_fma_rn_ftz_relu_f16x2(
-// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.rn.ftz.relu.f16x2(<2 x half> %{{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}})
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.ftz.relu.v2f16(<2 x half> %{{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}}, /* rnd=rn */ i32 1)
extern "C" __device__ f16x2 test_fma_rn_ftz_relu_f16x2(f16x2 a0, f16x2 a1, f16x2 a2) {
return __nvvm_fma_rn_ftz_relu_f16x2(a0, a1, a2);
}
// CIR-LABEL: @test_fma_rn_ftz_sat_f16
-// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.ftz.sat.f16" {{.*}} : (!cir.f16, !cir.f16, !cir.f16) -> !cir.f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.ftz.sat" {{.*}} : (!cir.f16, !cir.f16, !cir.f16, !s32i) -> !cir.f16
// LLVM-LABEL: @test_fma_rn_ftz_sat_f16(
-// LLVM: call{{.*}} half @llvm.nvvm.fma.rn.ftz.sat.f16(half %{{.*}}, half %{{.*}}, half %{{.*}})
+// LLVM: call{{.*}} half @llvm.nvvm.fma.ftz.sat.f16(half %{{.*}}, half %{{.*}}, half %{{.*}}, /* rnd=rn */ i32 1)
extern "C" __device__ _Float16 test_fma_rn_ftz_sat_f16(_Float16 a0, _Float16 a1, _Float16 a2) {
return __nvvm_fma_rn_ftz_sat_f16(a0, a1, a2);
}
// CIR-LABEL: @test_fma_rn_ftz_sat_f16x2
-// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.ftz.sat.f16x2" {{.*}} : (!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x !cir.f16>
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.ftz.sat" {{.*}} : (!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !s32i) -> !cir.vector<2 x !cir.f16>
// LLVM-LABEL: @test_fma_rn_ftz_sat_f16x2(
-// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.rn.ftz.sat.f16x2(<2 x half> %{{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}})
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.ftz.sat.v2f16(<2 x half> %{{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}}, /* rnd=rn */ i32 1)
extern "C" __device__ f16x2 test_fma_rn_ftz_sat_f16x2(f16x2 a0, f16x2 a1, f16x2 a2) {
return __nvvm_fma_rn_ftz_sat_f16x2(a0, a1, a2);
}
// CIR-LABEL: @test_fma_rn_relu_f16
-// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.relu.f16" {{.*}} : (!cir.f16, !cir.f16, !cir.f16) -> !cir.f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.relu" {{.*}} : (!cir.f16, !cir.f16, !cir.f16, !s32i) -> !cir.f16
// LLVM-LABEL: @test_fma_rn_relu_f16(
-// LLVM: call{{.*}} half @llvm.nvvm.fma.rn.relu.f16(half %{{.*}}, half %{{.*}}, half %{{.*}})
+// LLVM: call{{.*}} half @llvm.nvvm.fma.relu.f16(half %{{.*}}, half %{{.*}}, half %{{.*}}, /* rnd=rn */ i32 1)
extern "C" __device__ _Float16 test_fma_rn_relu_f16(_Float16 a0, _Float16 a1, _Float16 a2) {
return __nvvm_fma_rn_relu_f16(a0, a1, a2);
}
// CIR-LABEL: @test_fma_rn_relu_f16x2
-// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.relu.f16x2" {{.*}} : (!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>) -> !cir.vector<2 x !cir.f16>
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.relu" {{.*}} : (!cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !cir.vector<2 x !cir.f16>, !s32i) -> !cir.vector<2 x !cir.f16>
// LLVM-LABEL: @test_fma_rn_relu_f16x2(
-// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.rn.relu.f16x2(<2 x half> %{{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}})
+// LLVM: call{{.*}} <2 x half> @llvm.nvvm.fma.relu.v2f16(<2 x half> %{{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}}, /* rnd=rn */ i32 1)
extern "C" __device__ f16x2 test_fma_rn_relu_f16x2(f16x2 a0, f16x2 a1, f16x2 a2) {
return __nvvm_fma_rn_relu_f16x2(a0, a1, a2);
}
// CIR-LABEL: @test_fma_rn_sat_f16
-// CIR: cir.call_llvm_intrinsic "nvvm.fma.rn.sat.f16" {{.*}} : (!cir.f16, !cir.f16, !cir.f16) -> !cir.f16
+// CIR: cir.call_llvm_intrinsic "nvvm.fma.sat" {{.*}} : (!cir.f16, !cir.f16, !cir.f16, !s32i) -> !cir.f16
// LLVM-LABEL: @test_fma_rn_sat_f16(
-// LLVM: call{{.*}} half @llvm.nvvm.fma.rn.sat.f16(half %{{.*}}, half %{{.*}}, half %{{.*}})
+// LLVM: call{{.*}} half @llvm.nvvm.fma.sat.f16(half %{{.*}}, half %{{.*}}, half %{{.*}}, /* rnd=rn */ i32 1)
extern "C" __device__ _Float16 test_fma_rn_sat_f16(_Float16 a0, _Float16 a1, _Float16 a2) {
return __nvvm_fma_rn_sat_f16(a0, a1, a2);
}
// CIR-LABEL: @test_fma_rn_sat_f16x2
-// CIR: cir...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/232672
More information about the Mlir-commits
mailing list