[clang] [llvm] [clang][NVPTX] Add overloaded fmul intrinsics (PR #224546)
Srinivasa Ravi via llvm-commits
llvm-commits at lists.llvm.org
Tue Sep 22 00:32:46 PDT 2026
https://github.com/Wolfram70 updated https://github.com/llvm/llvm-project/pull/224546
>From 68067d172438e74341120b452de88d246f573091 Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Thu, 17 Sep 2026 11:59:07 +0000
Subject: [PATCH 1/2] [clang][NVPTX] Add overloaded fmul intrinsics
This change adds the following overloaded `fmul` intrinsics with
NVPTX codegen:
- `llvm.nvvm.fmul`
- `llvm.nvvm.fmul.ftz`
- `llvm.nvvm.fmul.sat`
- `llvm.nvvm.fmul.ftz.sat`
The rounding mode is passed in as an `i32` immediate operand.
Auto-upgrades the older non-overloaded intrinsics to the new
ones, and updates clang builtins and CIR codegen to lower to the
new intrinsics.
In the interest of completion, this also:
- Adds intrinsics support for lowering some multiplications that
were omitted earlier (`f16/f16x2` without saturation, `bf16/bf16x2`
multiplications, and `f32` with saturation), and support for the
`f32x2` type.
- Adds tests for constant folding of these intrinsics with the newly
supported scalar types.
PTX Spec Reference:
https://docs.nvidia.com/cuda/developer-preview/13.4/parallel-thread-execution/index.html#floating-point-instructions-mul
Assisted-by: Claude Opus 5
---
clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp | 109 ++-
clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp | 106 ++-
.../CIR/CodeGenCUDA/builtins-nvvm-math.cu | 24 +
clang/test/CodeGen/builtins-nvptx.c | 8 +-
llvm/docs/NVPTXUsage.md | 56 +-
llvm/include/llvm/IR/IntrinsicsNVVM.td | 30 +-
llvm/include/llvm/IR/NVVMIntrinsicUtils.h | 55 +-
llvm/lib/Analysis/ConstantFolding.cpp | 66 +-
llvm/lib/IR/AutoUpgrade.cpp | 53 +-
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 42 +-
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 10 +
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 66 +-
.../Assembler/auto_upgrade_nvvm_intrinsics.ll | 24 +
llvm/test/CodeGen/NVPTX/bf16-mul.ll | 33 +
llvm/test/CodeGen/NVPTX/f16-mul-sat.ll | 63 --
llvm/test/CodeGen/NVPTX/f16-mul.ll | 123 +++
llvm/test/CodeGen/NVPTX/fp-arith-sat.ll | 33 +
llvm/test/CodeGen/NVPTX/fp-mul-f32x2.ll | 125 +++
llvm/test/CodeGen/NVPTX/fp-mul-invalid.ll | 47 +
llvm/test/CodeGen/NVPTX/fp-mul.ll | 58 ++
.../InstCombine/NVPTX/nvvm-intrins.ll | 18 +-
.../InstSimplify/const-fold-nvvm-mul.ll | 832 +++++++++++++++---
llvm/test/Verifier/NVPTX/fmul.ll | 16 +
llvm/test/Verifier/intrinsic-bad-arg-type1.ll | 4 +-
llvm/unittests/IR/IntrinsicsTest.cpp | 2 +-
25 files changed, 1559 insertions(+), 444 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/bf16-mul.ll
delete mode 100644 llvm/test/CodeGen/NVPTX/f16-mul-sat.ll
create mode 100644 llvm/test/CodeGen/NVPTX/f16-mul.ll
create mode 100644 llvm/test/CodeGen/NVPTX/fp-mul-f32x2.ll
create mode 100644 llvm/test/CodeGen/NVPTX/fp-mul-invalid.ll
create mode 100644 llvm/test/CodeGen/NVPTX/fp-mul.ll
create mode 100644 llvm/test/Verifier/NVPTX/fmul.ll
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
index b81faf65414c1d..31bd0e820d7f36 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp
@@ -69,11 +69,11 @@ static mlir::Value emitUnaryNVVMIntrinsic(CIRGenFunction &cgf,
.getResult();
}
-/// Emit a CIR LLVMIntrinsicCallOp for an NVVM fadd intrinsic, which takes the
-/// rounding mode as a trailing operand.
-static mlir::Value emitNVVMFAdd(CIRGenFunction &cgf, const CallExpr *expr,
- llvm::StringRef intrinsicName,
- llvm::APFloat::roundingMode rm) {
+/// Emit a CIR LLVMIntrinsicCallOp for an NVVM fadd/fmul intrinsic, which takes
+/// the rounding mode as a trailing operand.
+static mlir::Value emitNVVMFPArith(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 lhs = cgf.emitScalarExpr(expr->getArg(0));
@@ -735,59 +735,96 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) {
return emitUnaryNVVMIntrinsic(*this, expr, "nvvm.ex2.approx.ftz");
case NVPTX::BI__nvvm_add_rn_f:
case NVPTX::BI__nvvm_add_rn_d:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd",
- llvm::APFloat::rmNearestTiesToEven);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd",
+ llvm::APFloat::rmNearestTiesToEven);
case NVPTX::BI__nvvm_add_rz_f:
case NVPTX::BI__nvvm_add_rz_d:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd", llvm::APFloat::rmTowardZero);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd",
+ llvm::APFloat::rmTowardZero);
case NVPTX::BI__nvvm_add_rm_f:
case NVPTX::BI__nvvm_add_rm_d:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd",
- llvm::APFloat::rmTowardNegative);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd",
+ llvm::APFloat::rmTowardNegative);
case NVPTX::BI__nvvm_add_rp_f:
case NVPTX::BI__nvvm_add_rp_d:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd",
- llvm::APFloat::rmTowardPositive);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd",
+ llvm::APFloat::rmTowardPositive);
case NVPTX::BI__nvvm_add_rn_ftz_f:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd.ftz",
- llvm::APFloat::rmNearestTiesToEven);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd.ftz",
+ llvm::APFloat::rmNearestTiesToEven);
case NVPTX::BI__nvvm_add_rz_ftz_f:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd.ftz",
- llvm::APFloat::rmTowardZero);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd.ftz",
+ llvm::APFloat::rmTowardZero);
case NVPTX::BI__nvvm_add_rm_ftz_f:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd.ftz",
- llvm::APFloat::rmTowardNegative);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd.ftz",
+ llvm::APFloat::rmTowardNegative);
case NVPTX::BI__nvvm_add_rp_ftz_f:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd.ftz",
- llvm::APFloat::rmTowardPositive);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd.ftz",
+ llvm::APFloat::rmTowardPositive);
case NVPTX::BI__nvvm_add_rn_sat_f:
case NVPTX::BI__nvvm_add_rn_sat_f16:
case NVPTX::BI__nvvm_add_rn_sat_v2f16:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd.sat",
- llvm::APFloat::rmNearestTiesToEven);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd.sat",
+ llvm::APFloat::rmNearestTiesToEven);
case NVPTX::BI__nvvm_add_rz_sat_f:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd.sat",
- llvm::APFloat::rmTowardZero);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd.sat",
+ llvm::APFloat::rmTowardZero);
case NVPTX::BI__nvvm_add_rm_sat_f:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd.sat",
- llvm::APFloat::rmTowardNegative);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd.sat",
+ llvm::APFloat::rmTowardNegative);
case NVPTX::BI__nvvm_add_rp_sat_f:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd.sat",
- llvm::APFloat::rmTowardPositive);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd.sat",
+ llvm::APFloat::rmTowardPositive);
case NVPTX::BI__nvvm_add_rn_ftz_sat_f:
case NVPTX::BI__nvvm_add_rn_ftz_sat_f16:
case NVPTX::BI__nvvm_add_rn_ftz_sat_v2f16:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd.ftz.sat",
- llvm::APFloat::rmNearestTiesToEven);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd.ftz.sat",
+ llvm::APFloat::rmNearestTiesToEven);
case NVPTX::BI__nvvm_add_rz_ftz_sat_f:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd.ftz.sat",
- llvm::APFloat::rmTowardZero);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd.ftz.sat",
+ llvm::APFloat::rmTowardZero);
case NVPTX::BI__nvvm_add_rm_ftz_sat_f:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd.ftz.sat",
- llvm::APFloat::rmTowardNegative);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd.ftz.sat",
+ llvm::APFloat::rmTowardNegative);
case NVPTX::BI__nvvm_add_rp_ftz_sat_f:
- return emitNVVMFAdd(*this, expr, "nvvm.fadd.ftz.sat",
- llvm::APFloat::rmTowardPositive);
+ return emitNVVMFPArith(*this, expr, "nvvm.fadd.ftz.sat",
+ llvm::APFloat::rmTowardPositive);
+ case NVPTX::BI__nvvm_mul_rn_f:
+ case NVPTX::BI__nvvm_mul_rn_d:
+ return emitNVVMFPArith(*this, expr, "nvvm.fmul",
+ llvm::APFloat::rmNearestTiesToEven);
+ case NVPTX::BI__nvvm_mul_rz_f:
+ case NVPTX::BI__nvvm_mul_rz_d:
+ return emitNVVMFPArith(*this, expr, "nvvm.fmul",
+ llvm::APFloat::rmTowardZero);
+ case NVPTX::BI__nvvm_mul_rm_f:
+ case NVPTX::BI__nvvm_mul_rm_d:
+ return emitNVVMFPArith(*this, expr, "nvvm.fmul",
+ llvm::APFloat::rmTowardNegative);
+ case NVPTX::BI__nvvm_mul_rp_f:
+ case NVPTX::BI__nvvm_mul_rp_d:
+ return emitNVVMFPArith(*this, expr, "nvvm.fmul",
+ llvm::APFloat::rmTowardPositive);
+ case NVPTX::BI__nvvm_mul_rn_ftz_f:
+ return emitNVVMFPArith(*this, expr, "nvvm.fmul.ftz",
+ llvm::APFloat::rmNearestTiesToEven);
+ case NVPTX::BI__nvvm_mul_rz_ftz_f:
+ return emitNVVMFPArith(*this, expr, "nvvm.fmul.ftz",
+ llvm::APFloat::rmTowardZero);
+ case NVPTX::BI__nvvm_mul_rm_ftz_f:
+ return emitNVVMFPArith(*this, expr, "nvvm.fmul.ftz",
+ llvm::APFloat::rmTowardNegative);
+ case NVPTX::BI__nvvm_mul_rp_ftz_f:
+ return emitNVVMFPArith(*this, expr, "nvvm.fmul.ftz",
+ llvm::APFloat::rmTowardPositive);
+ case NVPTX::BI__nvvm_mul_rn_sat_f16:
+ case NVPTX::BI__nvvm_mul_rn_sat_v2f16:
+ return emitNVVMFPArith(*this, expr, "nvvm.fmul.sat",
+ llvm::APFloat::rmNearestTiesToEven);
+ case NVPTX::BI__nvvm_mul_rn_ftz_sat_f16:
+ case NVPTX::BI__nvvm_mul_rn_ftz_sat_v2f16:
+ return emitNVVMFPArith(*this, expr, "nvvm.fmul.ftz.sat",
+ llvm::APFloat::rmNearestTiesToEven);
case NVPTX::BI__nvvm_ldg_h:
case NVPTX::BI__nvvm_ldg_h2:
cgm.errorNYI(expr->getSourceRange(),
diff --git a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
index 76f6757326ecab..830d2e2bb8f431 100644
--- a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
@@ -433,9 +433,9 @@ static Value *MakeFMAOOB(unsigned IntrinsicID, llvm::Type *Ty,
CGF.EmitScalarExpr(E->getArg(2))});
}
-static Value *MakeFAdd(unsigned IntrinsicID, APFloat::roundingMode RM,
- unsigned BuiltinID, const CallExpr *E,
- CodeGenFunction &CGF) {
+static Value *MakeFPArith(unsigned IntrinsicID, APFloat::roundingMode RM,
+ unsigned BuiltinID, const CallExpr *E,
+ CodeGenFunction &CGF) {
llvm::Type *Ty = CGF.ConvertType(E->getType());
return MakeHalfType(CGF.CGM.getIntrinsic(IntrinsicID, Ty), BuiltinID, E, CGF,
{CGF.Builder.getInt32(static_cast<int>(RM))});
@@ -1241,60 +1241,96 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
EmitScalarExpr(E->getArg(0)));
case NVPTX::BI__nvvm_add_rn_f:
case NVPTX::BI__nvvm_add_rn_d:
- return MakeFAdd(Intrinsic::nvvm_fadd, APFloat::rmNearestTiesToEven,
- BuiltinID, E, *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd, APFloat::rmNearestTiesToEven,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_add_rz_f:
case NVPTX::BI__nvvm_add_rz_d:
- return MakeFAdd(Intrinsic::nvvm_fadd, APFloat::rmTowardZero, BuiltinID, E,
- *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd, APFloat::rmTowardZero, BuiltinID,
+ E, *this);
case NVPTX::BI__nvvm_add_rm_f:
case NVPTX::BI__nvvm_add_rm_d:
- return MakeFAdd(Intrinsic::nvvm_fadd, APFloat::rmTowardNegative, BuiltinID,
- E, *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd, APFloat::rmTowardNegative,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_add_rp_f:
case NVPTX::BI__nvvm_add_rp_d:
- return MakeFAdd(Intrinsic::nvvm_fadd, APFloat::rmTowardPositive, BuiltinID,
- E, *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd, APFloat::rmTowardPositive,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_add_rn_ftz_f:
- return MakeFAdd(Intrinsic::nvvm_fadd_ftz, APFloat::rmNearestTiesToEven,
- BuiltinID, E, *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd_ftz, APFloat::rmNearestTiesToEven,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_add_rz_ftz_f:
- return MakeFAdd(Intrinsic::nvvm_fadd_ftz, APFloat::rmTowardZero, BuiltinID,
- E, *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd_ftz, APFloat::rmTowardZero,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_add_rm_ftz_f:
- return MakeFAdd(Intrinsic::nvvm_fadd_ftz, APFloat::rmTowardNegative,
- BuiltinID, E, *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd_ftz, APFloat::rmTowardNegative,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_add_rp_ftz_f:
- return MakeFAdd(Intrinsic::nvvm_fadd_ftz, APFloat::rmTowardPositive,
- BuiltinID, E, *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd_ftz, APFloat::rmTowardPositive,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_add_rn_sat_f:
case NVPTX::BI__nvvm_add_rn_sat_f16:
case NVPTX::BI__nvvm_add_rn_sat_v2f16:
- return MakeFAdd(Intrinsic::nvvm_fadd_sat, APFloat::rmNearestTiesToEven,
- BuiltinID, E, *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd_sat, APFloat::rmNearestTiesToEven,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_add_rz_sat_f:
- return MakeFAdd(Intrinsic::nvvm_fadd_sat, APFloat::rmTowardZero, BuiltinID,
- E, *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd_sat, APFloat::rmTowardZero,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_add_rm_sat_f:
- return MakeFAdd(Intrinsic::nvvm_fadd_sat, APFloat::rmTowardNegative,
- BuiltinID, E, *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd_sat, APFloat::rmTowardNegative,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_add_rp_sat_f:
- return MakeFAdd(Intrinsic::nvvm_fadd_sat, APFloat::rmTowardPositive,
- BuiltinID, E, *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd_sat, APFloat::rmTowardPositive,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_add_rn_ftz_sat_f:
case NVPTX::BI__nvvm_add_rn_ftz_sat_f16:
case NVPTX::BI__nvvm_add_rn_ftz_sat_v2f16:
- return MakeFAdd(Intrinsic::nvvm_fadd_ftz_sat, APFloat::rmNearestTiesToEven,
- BuiltinID, E, *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd_ftz_sat,
+ APFloat::rmNearestTiesToEven, BuiltinID, E, *this);
case NVPTX::BI__nvvm_add_rz_ftz_sat_f:
- return MakeFAdd(Intrinsic::nvvm_fadd_ftz_sat, APFloat::rmTowardZero,
- BuiltinID, E, *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd_ftz_sat, APFloat::rmTowardZero,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_add_rm_ftz_sat_f:
- return MakeFAdd(Intrinsic::nvvm_fadd_ftz_sat, APFloat::rmTowardNegative,
- BuiltinID, E, *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd_ftz_sat, APFloat::rmTowardNegative,
+ BuiltinID, E, *this);
case NVPTX::BI__nvvm_add_rp_ftz_sat_f:
- return MakeFAdd(Intrinsic::nvvm_fadd_ftz_sat, APFloat::rmTowardPositive,
- BuiltinID, E, *this);
+ return MakeFPArith(Intrinsic::nvvm_fadd_ftz_sat, APFloat::rmTowardPositive,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_mul_rn_f:
+ case NVPTX::BI__nvvm_mul_rn_d:
+ return MakeFPArith(Intrinsic::nvvm_fmul, APFloat::rmNearestTiesToEven,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_mul_rz_f:
+ case NVPTX::BI__nvvm_mul_rz_d:
+ return MakeFPArith(Intrinsic::nvvm_fmul, APFloat::rmTowardZero, BuiltinID,
+ E, *this);
+ case NVPTX::BI__nvvm_mul_rm_f:
+ case NVPTX::BI__nvvm_mul_rm_d:
+ return MakeFPArith(Intrinsic::nvvm_fmul, APFloat::rmTowardNegative,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_mul_rp_f:
+ case NVPTX::BI__nvvm_mul_rp_d:
+ return MakeFPArith(Intrinsic::nvvm_fmul, APFloat::rmTowardPositive,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_mul_rn_ftz_f:
+ return MakeFPArith(Intrinsic::nvvm_fmul_ftz, APFloat::rmNearestTiesToEven,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_mul_rz_ftz_f:
+ return MakeFPArith(Intrinsic::nvvm_fmul_ftz, APFloat::rmTowardZero,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_mul_rm_ftz_f:
+ return MakeFPArith(Intrinsic::nvvm_fmul_ftz, APFloat::rmTowardNegative,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_mul_rp_ftz_f:
+ return MakeFPArith(Intrinsic::nvvm_fmul_ftz, APFloat::rmTowardPositive,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_mul_rn_sat_f16:
+ case NVPTX::BI__nvvm_mul_rn_sat_v2f16:
+ return MakeFPArith(Intrinsic::nvvm_fmul_sat, APFloat::rmNearestTiesToEven,
+ BuiltinID, E, *this);
+ case NVPTX::BI__nvvm_mul_rn_ftz_sat_f16:
+ case NVPTX::BI__nvvm_mul_rn_ftz_sat_v2f16:
+ return MakeFPArith(Intrinsic::nvvm_fmul_ftz_sat,
+ APFloat::rmNearestTiesToEven, BuiltinID, E, *this);
case NVPTX::BI__nvvm_ldg_h:
case NVPTX::BI__nvvm_ldg_h2:
return MakeLdg(*this, E);
diff --git a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-math.cu b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-math.cu
index 69df1d376f7f73..9660598f680d77 100644
--- a/clang/test/CIR/CodeGenCUDA/builtins-nvvm-math.cu
+++ b/clang/test/CIR/CodeGenCUDA/builtins-nvvm-math.cu
@@ -87,3 +87,27 @@ __device__ double test_add_rz_d(double x, double y) {
__device__ float test_add_rm_ftz_sat_f(float x, float y) {
return __nvvm_add_rm_ftz_sat_f(x, y);
}
+
+// CIR-LABEL: @_Z13test_mul_rn_fff
+// CIR: cir.call_llvm_intrinsic "nvvm.fmul" {{.*}} : (!cir.float, !cir.float, !s32i) -> !cir.float
+// LLVM-LABEL: @_Z13test_mul_rn_fff
+// LLVM: call {{.*}}float @llvm.nvvm.fmul.f32(float {{.*}}, float {{.*}}, /* rnd=rn */ i32 1)
+__device__ float test_mul_rn_f(float x, float y) {
+ return __nvvm_mul_rn_f(x, y);
+}
+
+// CIR-LABEL: @_Z13test_mul_rz_ddd
+// CIR: cir.call_llvm_intrinsic "nvvm.fmul" {{.*}} : (!cir.double, !cir.double, !s32i) -> !cir.double
+// LLVM-LABEL: @_Z13test_mul_rz_ddd
+// LLVM: call {{.*}}double @llvm.nvvm.fmul.f64(double {{.*}}, double {{.*}}, /* rnd=rz */ i32 0)
+__device__ double test_mul_rz_d(double x, double y) {
+ return __nvvm_mul_rz_d(x, y);
+}
+
+// CIR-LABEL: @_Z17test_mul_rp_ftz_fff
+// CIR: cir.call_llvm_intrinsic "nvvm.fmul.ftz" {{.*}} : (!cir.float, !cir.float, !s32i) -> !cir.float
+// LLVM-LABEL: @_Z17test_mul_rp_ftz_fff
+// LLVM: call {{.*}}float @llvm.nvvm.fmul.ftz.f32(float {{.*}}, float {{.*}}, /* rnd=rp */ i32 2)
+__device__ float test_mul_rp_ftz_f(float x, float y) {
+ return __nvvm_mul_rp_ftz_f(x, y);
+}
diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c
index 607095bce481c7..a6c1a9c3e72725 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -1700,13 +1700,13 @@ __device__ void nvvm_add_mul_f16_sat() {
// CHECK: call <2 x half> @llvm.nvvm.fadd.ftz.sat.v2f16({{.*}}i32 1)
__nvvm_add_rn_ftz_sat_v2f16(F16X2, F16X2_2);
- // CHECK: call half @llvm.nvvm.mul.rn.sat.f16
+ // CHECK: call half @llvm.nvvm.fmul.sat.f16({{.*}}i32 1)
__nvvm_mul_rn_sat_f16(F16, F16_2);
- // CHECK: call half @llvm.nvvm.mul.rn.ftz.sat.f16
+ // CHECK: call half @llvm.nvvm.fmul.ftz.sat.f16({{.*}}i32 1)
__nvvm_mul_rn_ftz_sat_f16(F16, F16_2);
- // CHECK: call <2 x half> @llvm.nvvm.mul.rn.sat.v2f16
+ // CHECK: call <2 x half> @llvm.nvvm.fmul.sat.v2f16({{.*}}i32 1)
__nvvm_mul_rn_sat_v2f16(F16X2, F16X2_2);
- // CHECK: call <2 x half> @llvm.nvvm.mul.rn.ftz.sat.v2f16
+ // CHECK: call <2 x half> @llvm.nvvm.fmul.ftz.sat.v2f16({{.*}}i32 1)
__nvvm_mul_rn_ftz_sat_v2f16(F16X2, F16X2_2);
// CHECK: ret void
diff --git a/llvm/docs/NVPTXUsage.md b/llvm/docs/NVPTXUsage.md
index 67bf8ea482933a..1386046f7b85d2 100644
--- a/llvm/docs/NVPTXUsage.md
+++ b/llvm/docs/NVPTXUsage.md
@@ -1403,29 +1403,65 @@ PTX instruction. The supported combinations are:
- None
```
-#### '`llvm.nvvm.mul.*`' Half-precision Intrinsics
+#### '`llvm.nvvm.fmul.*`' Intrinsics
##### Syntax:
-```llvm
-declare half @llvm.nvvm.mul.rn.sat.f16(half %a, half %b)
-declare <2 x half> @llvm.nvvm.mul.rn.sat.v2f16(<2 x half> %a, <2 x half> %b)
+This is an overloaded intrinsic. The '`.ftz`' and '`.sat`' modifiers are
+optional.
-declare half @llvm.nvvm.mul.rn.ftz.sat.f16(half %a, half %b)
-declare <2 x half> @llvm.nvvm.mul.rn.ftz.sat.v2f16(<2 x half> %a, <2 x half> %b)
+```llvm
+declare half @llvm.nvvm.fmul{.ftz}{.sat}.f16(half %a, half %b, i32 immarg %rnd)
+declare <2 x half> @llvm.nvvm.fmul{.ftz}{.sat}.v2f16(<2 x half> %a, <2 x half> %b, i32 immarg %rnd)
+declare bfloat @llvm.nvvm.fmul.bf16(bfloat %a, bfloat %b, i32 immarg %rnd)
+declare <2 x bfloat> @llvm.nvvm.fmul.v2bf16(<2 x bfloat> %a, <2 x bfloat> %b, i32 immarg %rnd)
+declare float @llvm.nvvm.fmul{.ftz}{.sat}.f32(float %a, float %b, i32 immarg %rnd)
+declare <2 x float> @llvm.nvvm.fmul{.ftz}.v2f32(<2 x float> %a, <2 x float> %b, i32 immarg %rnd)
+declare double @llvm.nvvm.fmul.f64(double %a, double %b, i32 immarg %rnd)
```
##### Overview:
-The '`llvm.nvvm.mul.*`' intrinsics perform a multiplication operation with
-the specified rounding mode and modifiers.
+The '`llvm.nvvm.fmul.*`' intrinsics multiply `%a` and `%b` using the rounding
+mode selected by `%rnd` and the modifiers present in the intrinsic name. They
+correspond directly to the `mul` PTX instruction.
##### Semantics:
-The '`.sat`' modifier performs a saturating multiplication where the result is
-clamped to `[0.0, 1.0]` and `NaN` results are flushed to `+0.0f`.
+`%rnd` selects the rounding mode applied to the result, see
+{ref}`fp-rounding-modes`.
+
The '`.ftz`' modifier flushes subnormal inputs and results to sign-preserving
zero.
+The '`.sat`' modifier performs a saturating multiplication where the result is
+clamped to `[0.0, 1.0]` and `NaN` results are flushed to `+0.0f`.
+
+Not every combination of operand type, rounding mode and modifier maps to a
+PTX instruction. The supported combinations are:
+
+```{list-table}
+:widths: 25 25 25
+:header-rows: 1
+
+ * - Operand Type
+ - Rounding Modes
+ - Modifiers
+ * - `half`, `<2 x half>`
+ - `rn`
+ - `.ftz`, `.sat`
+ * - `bfloat`, `<2 x bfloat>`
+ - `rn`
+ - None
+ * - `float`
+ - `rn`, `rz`, `rp`, `rm`
+ - `.ftz`, `.sat`
+ * - `<2 x float>`
+ - `rn`, `rz`, `rp`, `rm`
+ - `.ftz`
+ * - `double`
+ - `rn`, `rz`, `rp`, `rm`
+ - None
+```
#### '`llvm.nvvm.fma.*`' Half-precision Intrinsics
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index ff12175180d149..50d32c305571b6 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1678,25 +1678,21 @@ let TargetPrefix = "nvvm" in {
def int_nvvm_mul24_ # sign # i : NVVMBuiltin,
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty]>;
}
-
- foreach rnd = ["rn", "rz", "rm", "rp"] in {
- foreach ftz = ["", "_ftz"] in
- def int_nvvm_mul_ # rnd # ftz # _f : NVVMBuiltin,
- DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty]>;
-
- def int_nvvm_mul_ # rnd # _d : NVVMBuiltin,
- DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty]>;
- }
-
- foreach ftz = ["", "_ftz"] in {
- def int_nvvm_mul_rn # ftz # _sat_f16 : NVVMBuiltin,
- DefaultAttrsIntrinsic<[llvm_half_ty], [llvm_half_ty, llvm_half_ty]>;
-
- def int_nvvm_mul_rn # ftz # _sat_v2f16 : NVVMBuiltin,
- DefaultAttrsIntrinsic<[llvm_v2f16_ty], [llvm_v2f16_ty, llvm_v2f16_ty]>;
- } // ftz
}
+ let IntrProperties = [IntrNoMem, IntrSpeculatable, Commutative,
+ IntrNoCreateUndefOrPoison, ImmArg<ArgIndex<2>>,
+ Range<ArgIndex<2>, 0, 4>,
+ ArgInfo<ArgIndex<2>,
+ [ArgName<"rnd">,
+ ImmArgPrinter<"printFPRoundingMode">]>] in
+ foreach ftz = ["", "_ftz"] in
+ foreach sat = ["", "_sat"] in
+ def int_nvvm_fmul # ftz # sat :
+ DefaultAttrsIntrinsic<[llvm_anyfloat_ty],
+ [LLVMMatchType<0>, LLVMMatchType<0>,
+ llvm_i32_ty]>;
+
//
// Div
//
diff --git a/llvm/include/llvm/IR/NVVMIntrinsicUtils.h b/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
index 9585cdac2a4ea9..40efa898d27281 100644
--- a/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
+++ b/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
@@ -697,47 +697,38 @@ inline StringRef GetRoundingModeName(APFloat::roundingMode RM) {
}
}
-inline bool FMulShouldFTZ(Intrinsic::ID IntrinsicID) {
+inline bool FPArithShouldFTZ(Intrinsic::ID IntrinsicID) {
switch (IntrinsicID) {
- case Intrinsic::nvvm_mul_rm_ftz_f:
- case Intrinsic::nvvm_mul_rn_ftz_f:
- case Intrinsic::nvvm_mul_rp_ftz_f:
- case Intrinsic::nvvm_mul_rz_ftz_f:
+ case Intrinsic::nvvm_fadd_ftz:
+ case Intrinsic::nvvm_fadd_ftz_sat:
+ case Intrinsic::nvvm_fmul_ftz:
+ case Intrinsic::nvvm_fmul_ftz_sat:
return true;
- case Intrinsic::nvvm_mul_rm_f:
- case Intrinsic::nvvm_mul_rn_f:
- case Intrinsic::nvvm_mul_rp_f:
- case Intrinsic::nvvm_mul_rz_f:
- case Intrinsic::nvvm_mul_rm_d:
- case Intrinsic::nvvm_mul_rn_d:
- case Intrinsic::nvvm_mul_rp_d:
- case Intrinsic::nvvm_mul_rz_d:
+ case Intrinsic::nvvm_fadd:
+ case Intrinsic::nvvm_fadd_sat:
+ case Intrinsic::nvvm_fmul:
+ case Intrinsic::nvvm_fmul_sat:
return false;
}
- llvm_unreachable("Checking FTZ flag for invalid NVVM mul intrinsic");
+ llvm_unreachable("Checking FTZ flag for invalid NVVM fadd/fmul intrinsic");
}
-inline APFloat::roundingMode GetFMulRoundingMode(Intrinsic::ID IntrinsicID) {
+inline bool FPArithIsSat(Intrinsic::ID IntrinsicID) {
switch (IntrinsicID) {
- case Intrinsic::nvvm_mul_rm_f:
- case Intrinsic::nvvm_mul_rm_d:
- case Intrinsic::nvvm_mul_rm_ftz_f:
- return APFloat::rmTowardNegative;
- case Intrinsic::nvvm_mul_rn_f:
- case Intrinsic::nvvm_mul_rn_d:
- case Intrinsic::nvvm_mul_rn_ftz_f:
- return APFloat::rmNearestTiesToEven;
- case Intrinsic::nvvm_mul_rp_f:
- case Intrinsic::nvvm_mul_rp_d:
- case Intrinsic::nvvm_mul_rp_ftz_f:
- return APFloat::rmTowardPositive;
- case Intrinsic::nvvm_mul_rz_f:
- case Intrinsic::nvvm_mul_rz_d:
- case Intrinsic::nvvm_mul_rz_ftz_f:
- return APFloat::rmTowardZero;
+ case Intrinsic::nvvm_fadd_sat:
+ case Intrinsic::nvvm_fadd_ftz_sat:
+ case Intrinsic::nvvm_fmul_sat:
+ case Intrinsic::nvvm_fmul_ftz_sat:
+ return true;
+
+ case Intrinsic::nvvm_fadd:
+ case Intrinsic::nvvm_fadd_ftz:
+ case Intrinsic::nvvm_fmul:
+ case Intrinsic::nvvm_fmul_ftz:
+ return false;
}
- llvm_unreachable("Invalid FP instrinsic rounding mode for NVVM mul");
+ llvm_unreachable("Checking sat flag for invalid NVVM fadd/fmul intrinsic");
}
inline bool FDivShouldFTZ(Intrinsic::ID IntrinsicID) {
diff --git a/llvm/lib/Analysis/ConstantFolding.cpp b/llvm/lib/Analysis/ConstantFolding.cpp
index a822ab9c5a748d..bca11a4dac0da6 100644
--- a/llvm/lib/Analysis/ConstantFolding.cpp
+++ b/llvm/lib/Analysis/ConstantFolding.cpp
@@ -1997,9 +1997,11 @@ static bool canConstantFoldIntrinsic(Intrinsic::ID ID, bool IsStrictFP) {
case Intrinsic::nvvm_sqrt_rn_ftz_f:
return !IsStrictFP;
- // NVVM add intrinsics with explicit rounding modes
+ // NVVM add/mul intrinsics with explicit rounding modes
case Intrinsic::nvvm_fadd:
case Intrinsic::nvvm_fadd_ftz:
+ case Intrinsic::nvvm_fmul:
+ case Intrinsic::nvvm_fmul_ftz:
// NVVM div intrinsics with explicit rounding modes
case Intrinsic::nvvm_div_rm_d:
@@ -2015,20 +2017,6 @@ static bool canConstantFoldIntrinsic(Intrinsic::ID ID, bool IsStrictFP) {
case Intrinsic::nvvm_div_rp_ftz_f:
case Intrinsic::nvvm_div_rz_ftz_f:
- // NVVM mul intrinsics with explicit rounding modes
- case Intrinsic::nvvm_mul_rm_d:
- case Intrinsic::nvvm_mul_rn_d:
- case Intrinsic::nvvm_mul_rp_d:
- case Intrinsic::nvvm_mul_rz_d:
- case Intrinsic::nvvm_mul_rm_f:
- case Intrinsic::nvvm_mul_rn_f:
- case Intrinsic::nvvm_mul_rp_f:
- case Intrinsic::nvvm_mul_rz_f:
- case Intrinsic::nvvm_mul_rm_ftz_f:
- case Intrinsic::nvvm_mul_rn_ftz_f:
- case Intrinsic::nvvm_mul_rp_ftz_f:
- case Intrinsic::nvvm_mul_rz_ftz_f:
-
// NVVM fma intrinsics with explicit rounding modes
case Intrinsic::nvvm_fma_rm_d:
case Intrinsic::nvvm_fma_rn_d:
@@ -3671,37 +3659,6 @@ static Constant *ConstantFoldIntrinsicCall2(Intrinsic::ID IntrinsicID, Type *Ty,
return ConstantFP::get(Ty, Res);
}
- case Intrinsic::nvvm_mul_rm_f:
- case Intrinsic::nvvm_mul_rn_f:
- case Intrinsic::nvvm_mul_rp_f:
- case Intrinsic::nvvm_mul_rz_f:
- case Intrinsic::nvvm_mul_rm_d:
- case Intrinsic::nvvm_mul_rn_d:
- case Intrinsic::nvvm_mul_rp_d:
- case Intrinsic::nvvm_mul_rz_d:
- case Intrinsic::nvvm_mul_rm_ftz_f:
- case Intrinsic::nvvm_mul_rn_ftz_f:
- case Intrinsic::nvvm_mul_rp_ftz_f:
- case Intrinsic::nvvm_mul_rz_ftz_f: {
-
- bool IsFTZ = nvvm::FMulShouldFTZ(IntrinsicID);
- APFloat A = IsFTZ ? FTZPreserveSign(Op1V) : Op1V;
- APFloat B = IsFTZ ? FTZPreserveSign(Op2V) : Op2V;
-
- APFloat::roundingMode RoundMode =
- nvvm::GetFMulRoundingMode(IntrinsicID);
-
- APFloat Res = A;
- APFloat::opStatus Status = Res.multiply(B, RoundMode);
-
- if (!Res.isNaN() &&
- (Status == APFloat::opOK || Status == APFloat::opInexact)) {
- Res = IsFTZ ? FTZPreserveSign(Res) : Res;
- return ConstantFP::get(Ty, Res);
- }
- return nullptr;
- }
-
case Intrinsic::nvvm_div_rm_f:
case Intrinsic::nvvm_div_rn_f:
case Intrinsic::nvvm_div_rp_f:
@@ -4229,17 +4186,24 @@ static Constant *ConstantFoldScalarCall3(StringRef Name,
}
// TODO: Add constant folding for the _sat variants.
- if (IntrinsicID == Intrinsic::nvvm_fadd ||
- IntrinsicID == Intrinsic::nvvm_fadd_ftz) {
- bool IsFTZ = IntrinsicID == Intrinsic::nvvm_fadd_ftz;
+ const bool IsFAdd = IntrinsicID == Intrinsic::nvvm_fadd ||
+ IntrinsicID == Intrinsic::nvvm_fadd_ftz;
+ const bool IsFMul = IntrinsicID == Intrinsic::nvvm_fmul ||
+ IntrinsicID == Intrinsic::nvvm_fmul_ftz;
+ if (IsFAdd || IsFMul) {
+ bool IsFTZ = IntrinsicID == Intrinsic::nvvm_fadd_ftz ||
+ IntrinsicID == Intrinsic::nvvm_fmul_ftz;
APFloat A =
IsFTZ ? FTZPreserveSign(Op1->getValueAPF()) : Op1->getValueAPF();
APFloat B =
IsFTZ ? FTZPreserveSign(Op2->getValueAPF()) : Op2->getValueAPF();
+ APFloat::roundingMode RoundMode =
+ nvvm::GetRoundingModeFromImmArg(Operands[2]);
+
APFloat Res = A;
APFloat::opStatus Status =
- Res.add(B, nvvm::GetRoundingModeFromImmArg(Operands[2]));
+ IsFAdd ? Res.add(B, RoundMode) : Res.multiply(B, RoundMode);
if (!Res.isNaN() &&
(Status == APFloat::opOK || Status == APFloat::opInexact)) {
@@ -4538,6 +4502,8 @@ static Constant *ConstantFoldFixedVectorCall(
}
case Intrinsic::nvvm_fadd:
case Intrinsic::nvvm_fadd_ftz:
+ case Intrinsic::nvvm_fmul:
+ case Intrinsic::nvvm_fmul_ftz:
// The rounding mode operand is a scalar, so the lane-wise folding below
// does not apply.
// TODO: Fold these by passing the rounding mode through to every lane.
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index 75fc03c22222bb..ad0ccdf6396c36 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -1505,8 +1505,16 @@ static bool isLegacyNVPTXBF16IntSignature(Function *F, Intrinsic::ID IID) {
return true;
}
+// Overloaded fadd/fmul intrinsic IDs, indexed by [`.ftz`][`.sat`].
+static constexpr Intrinsic::ID NVVMFAddIIDs[2][2] = {
+ {Intrinsic::nvvm_fadd, Intrinsic::nvvm_fadd_sat},
+ {Intrinsic::nvvm_fadd_ftz, Intrinsic::nvvm_fadd_ftz_sat}};
+static constexpr Intrinsic::ID NVVMFMulIIDs[2][2] = {
+ {Intrinsic::nvvm_fmul, Intrinsic::nvvm_fmul_sat},
+ {Intrinsic::nvvm_fmul_ftz, Intrinsic::nvvm_fmul_ftz_sat}};
+
static std::optional<std::pair<Intrinsic::ID, RoundingMode>>
-getNVVMFAddUpgrade(StringRef Name) {
+getNVVMFPArithUpgrade(StringRef Name, const Intrinsic::ID IIDs[2][2]) {
auto [Modifiers, Type] = Name.rsplit('.');
if (!is_contained({"f", "d", "f16", "v2f16"}, Type))
return std::nullopt;
@@ -1521,16 +1529,13 @@ getNVVMFAddUpgrade(StringRef Name) {
if (!RoundingMode)
return std::nullopt;
- Intrinsic::ID IID = StringSwitch<Intrinsic::ID>(Modifiers.drop_front(2))
- .Case("", Intrinsic::nvvm_fadd)
- .Case(".ftz", Intrinsic::nvvm_fadd_ftz)
- .Case(".sat", Intrinsic::nvvm_fadd_sat)
- .Case(".ftz.sat", Intrinsic::nvvm_fadd_ftz_sat)
- .Default(Intrinsic::not_intrinsic);
- if (IID == Intrinsic::not_intrinsic)
+ StringRef Rest = Modifiers.drop_front(2);
+ const bool IsFTZ = Rest.consume_front(".ftz");
+ const bool IsSat = Rest.consume_front(".sat");
+ if (!Rest.empty())
return std::nullopt;
- return std::make_pair(IID, *RoundingMode);
+ return std::make_pair(IIDs[IsFTZ][IsSat], *RoundingMode);
}
static Intrinsic::ID shouldUpgradeNVPTXMBarrierInitIntrinsic(StringRef Name) {
@@ -2166,7 +2171,10 @@ static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn,
Expand = Name == "f" || Name == "ftz.f" || Name == "d";
else if (Name.consume_front("add."))
// nvvm.add.<rnd>{.ftz}{.sat}.{f,d,f16,v2f16}
- Expand = getNVVMFAddUpgrade(Name).has_value();
+ Expand = getNVVMFPArithUpgrade(Name, NVVMFAddIIDs).has_value();
+ else if (Name.consume_front("mul."))
+ // nvvm.mul.<rnd>{.ftz}{.sat}.{f,d,f16,v2f16}
+ Expand = getNVVMFPArithUpgrade(Name, NVVMFMulIIDs).has_value();
else if (Name.consume_front("ex2.approx."))
// nvvm.ex2.approx.{f,ftz.f,d,f16x2}
Expand =
@@ -3227,6 +3235,19 @@ void llvm::UpgradeInlineAsmString(std::string *AsmStr) {
}
}
+static Value *upgradeNVVMFPArithCall(IRBuilder<> &Builder, CallBase *CI,
+ StringRef Name,
+ const Intrinsic::ID IIDs[2][2]) {
+ auto Upgrade = getNVVMFPArithUpgrade(Name, IIDs);
+ assert(Upgrade && "unsupported nvvm.add.*/nvvm.mul.* intrinsic");
+ auto [IID, RoundingMode] = *Upgrade;
+ Value *A = CI->getArgOperand(0);
+ return Builder.CreateIntrinsic(
+ A->getType(), IID,
+ {A, CI->getArgOperand(1),
+ Builder.getInt32(static_cast<int>(RoundingMode))});
+}
+
static Value *upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI,
Function *F, IRBuilder<> &Builder) {
Value *Rep = nullptr;
@@ -3249,14 +3270,10 @@ static Value *upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI,
Rep = Builder.CreateUnaryIntrinsic(IID, CI->getArgOperand(0));
} else if (Name.consume_front("add.")) {
// nvvm.add.<rnd>{.ftz}{.sat}.{f,d,f16,v2f16}
- auto FAdd = getNVVMFAddUpgrade(Name);
- assert(FAdd && "unsupported nvvm.add.* intrinsic");
- auto [IID, RoundingMode] = *FAdd;
- Value *A = CI->getArgOperand(0);
- Rep = Builder.CreateIntrinsic(
- A->getType(), IID,
- {A, CI->getArgOperand(1),
- Builder.getInt32(static_cast<int>(RoundingMode))});
+ Rep = upgradeNVVMFPArithCall(Builder, CI, Name, NVVMFAddIIDs);
+ } else if (Name.consume_front("mul.")) {
+ // nvvm.mul.<rnd>{.ftz}{.sat}.{f,d,f16,v2f16}
+ Rep = upgradeNVVMFPArithCall(Builder, CI, Name, NVVMFMulIIDs);
} else if (Name.consume_front("ex2.approx.")) {
// nvvm.ex2.approx.{f,ftz.f,d,f16x2}
Intrinsic::ID IID = Name.starts_with("ftz") ? Intrinsic::nvvm_ex2_approx_ftz
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index c01506f94ad90a..3ad43d336f1760 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -7178,10 +7178,8 @@ static SDValue sinkProxyReg(SDValue R, SDValue Chain,
static unsigned getFAddWithNegOpcode(EVT VT, Intrinsic::ID IID,
APFloat::roundingMode RoundingMode) {
- const bool IsFTZ =
- IID == Intrinsic::nvvm_fadd_ftz || IID == Intrinsic::nvvm_fadd_ftz_sat;
- const bool IsSat =
- IID == Intrinsic::nvvm_fadd_sat || IID == Intrinsic::nvvm_fadd_ftz_sat;
+ const bool IsFTZ = nvvm::FPArithShouldFTZ(IID);
+ const bool IsSat = nvvm::FPArithIsSat(IID);
switch (VT.getScalarType().getSimpleVT().SimpleTy) {
case MVT::f16: {
static constexpr unsigned SubRNOpcodes[2][2] = {
@@ -7236,22 +7234,20 @@ static SDValue combineFAddWithNeg(SDNode *N, SelectionDAG &DAG,
// TODO: Remove the type-legality checks here once
// https://github.com/llvm/llvm-project/pull/172442 lands, adding support for
// explicit type constraints for overloaded intrinsics in tablegen.
-static bool isSupportedFAdd(EVT VT, const NVPTXSubtarget &STI,
- Intrinsic::ID IID,
- APFloat::roundingMode RoundingMode) {
+static bool isSupportedFPArith(EVT VT, const NVPTXSubtarget &STI,
+ unsigned ISDOpcode, Intrinsic::ID IID,
+ APFloat::roundingMode RoundingMode) {
if (VT.isVector() && VT.getVectorElementCount() != ElementCount::getFixed(2))
return false;
const bool IsRN = RoundingMode == APFloat::rmNearestTiesToEven;
- const bool IsFTZ =
- IID == Intrinsic::nvvm_fadd_ftz || IID == Intrinsic::nvvm_fadd_ftz_sat;
- const bool IsSat =
- IID == Intrinsic::nvvm_fadd_sat || IID == Intrinsic::nvvm_fadd_ftz_sat;
+ const bool IsFTZ = nvvm::FPArithShouldFTZ(IID);
+ const bool IsSat = nvvm::FPArithIsSat(IID);
switch (VT.getScalarType().getSimpleVT().SimpleTy) {
case MVT::f16:
return IsRN;
case MVT::bf16:
- return IsRN && !IsSat && !IsFTZ && STI.hasNativeBF16Support(ISD::FADD);
+ return IsRN && !IsSat && !IsFTZ && STI.hasNativeBF16Support(ISDOpcode);
case MVT::f32:
return !VT.isVector() || (!IsSat && STI.hasF32x2Instructions());
case MVT::f64:
@@ -7261,9 +7257,9 @@ static bool isSupportedFAdd(EVT VT, const NVPTXSubtarget &STI,
}
}
-static SDValue diagnoseUnsupportedFAdd(SDNode *N, SelectionDAG &DAG,
- Intrinsic::ID IID,
- APFloat::roundingMode RoundingMode) {
+static SDValue diagnoseUnsupportedFPArith(SDNode *N, SelectionDAG &DAG,
+ Intrinsic::ID IID,
+ APFloat::roundingMode RoundingMode) {
const EVT VT = N->getValueType(0);
DAG.getContext()->diagnose(DiagnosticInfoUnsupported(
DAG.getMachineFunction().getFunction(),
@@ -7289,10 +7285,22 @@ static SDValue combineIntrinsicWOChain(SDNode *N,
case Intrinsic::nvvm_fadd_ftz_sat: {
const auto RoundingMode = static_cast<APFloat::roundingMode>(
N->getConstantOperandAPInt(3).getSExtValue());
- if (!isSupportedFAdd(N->getValueType(0), STI, IID, RoundingMode))
- return diagnoseUnsupportedFAdd(N, DCI.DAG, IID, RoundingMode);
+ if (!isSupportedFPArith(N->getValueType(0), STI, ISD::FADD, IID,
+ RoundingMode))
+ return diagnoseUnsupportedFPArith(N, DCI.DAG, IID, RoundingMode);
return combineFAddWithNeg(N, DCI.DAG, IID, RoundingMode);
}
+ case Intrinsic::nvvm_fmul:
+ case Intrinsic::nvvm_fmul_ftz:
+ case Intrinsic::nvvm_fmul_sat:
+ case Intrinsic::nvvm_fmul_ftz_sat: {
+ const auto RoundingMode = static_cast<APFloat::roundingMode>(
+ N->getConstantOperandAPInt(3).getSExtValue());
+ if (!isSupportedFPArith(N->getValueType(0), STI, ISD::FMUL, IID,
+ RoundingMode))
+ return diagnoseUnsupportedFPArith(N, DCI.DAG, IID, RoundingMode);
+ break;
+ }
}
return SDValue();
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index c130a2f351ff37..28ba2f05ee4471 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -174,6 +174,8 @@ def hasSHFL : PredNot<PredAnd<[SM70, PTX64]>>;
def allowFP16Math : SubtargetPredicate;
def hasBF16Math : SubtargetPredicate;
+defvar BF16ArithPreds = [hasBF16Math, PTX78, SM90];
+
def hasCLMAD : SubtargetPredicate;
def hasTcgen05InstSupport : SubtargetPredicate;
@@ -286,6 +288,14 @@ def fpimm_0 : FPImmLeaf<fAny, [{ return Imm.isZero(); }]>;
def fpimm_1 : FPImmLeaf<fAny, [{ return Imm.isOne(); }]>;
def fpimm_neg_1 : FPImmLeaf<fAny, [{ return Imm.isMinusOne(); }]>;
+class FPRndModeImm<string mode> : TImmLeaf<i32,
+ "return Imm == static_cast<int>(RoundingMode::" # mode # ");">;
+
+def fp_rnd_rn_imm : FPRndModeImm<"NearestTiesToEven">;
+def fp_rnd_rz_imm : FPRndModeImm<"TowardZero">;
+def fp_rnd_rm_imm : FPRndModeImm<"TowardNegative">;
+def fp_rnd_rp_imm : FPRndModeImm<"TowardPositive">;
+
// Operands which can hold a Register or an Immediate.
//
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index f54fd1b348af85..c245997e7c2f45 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2011,6 +2011,8 @@ class F_MATH_3<string OpcStr, NVPTXRegClass t_regclass,
(IntOP s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2))]>,
Requires<Preds>;
+defvar FPRoundingModes = ["rn", "rz", "rm", "rp"];
+
//
// MISC
//
@@ -2181,27 +2183,43 @@ def : Pat<(int_nvvm_mulhi_ui i32:$a, i32:$b), (MUL_HI_U32rr $a, $b)>;
def : Pat<(int_nvvm_mulhi_ll i64:$a, i64:$b), (MUL_HI_S64rr $a, $b)>;
def : Pat<(int_nvvm_mulhi_ull i64:$a, i64:$b), (MUL_HI_U64rr $a, $b)>;
-def INT_NVVM_MUL_RN_FTZ_F : F_MATH_2<"mul.rn.ftz.f32", B32, B32, B32, int_nvvm_mul_rn_ftz_f>;
-def INT_NVVM_MUL_RN_F : F_MATH_2<"mul.rn.f32", B32, B32, B32, int_nvvm_mul_rn_f>;
-def INT_NVVM_MUL_RZ_FTZ_F : F_MATH_2<"mul.rz.ftz.f32", B32, B32, B32, int_nvvm_mul_rz_ftz_f>;
-def INT_NVVM_MUL_RZ_F : F_MATH_2<"mul.rz.f32", B32, B32, B32, int_nvvm_mul_rz_f>;
-def INT_NVVM_MUL_RM_FTZ_F : F_MATH_2<"mul.rm.ftz.f32", B32, B32, B32, int_nvvm_mul_rm_ftz_f>;
-def INT_NVVM_MUL_RM_F : F_MATH_2<"mul.rm.f32", B32, B32, B32, int_nvvm_mul_rm_f>;
-def INT_NVVM_MUL_RP_FTZ_F : F_MATH_2<"mul.rp.ftz.f32", B32, B32, B32, int_nvvm_mul_rp_ftz_f>;
-def INT_NVVM_MUL_RP_F : F_MATH_2<"mul.rp.f32", B32, B32, B32, int_nvvm_mul_rp_f>;
-
-def INT_NVVM_MUL_RN_D : F_MATH_2<"mul.rn.f64", B64, B64, B64, int_nvvm_mul_rn_d>;
-def INT_NVVM_MUL_RZ_D : F_MATH_2<"mul.rz.f64", B64, B64, B64, int_nvvm_mul_rz_d>;
-def INT_NVVM_MUL_RM_D : F_MATH_2<"mul.rm.f64", B64, B64, B64, int_nvvm_mul_rm_d>;
-def INT_NVVM_MUL_RP_D : F_MATH_2<"mul.rp.f64", B64, B64, B64, int_nvvm_mul_rp_d>;
-
def INT_NVVM_MUL24_I : F_MATH_2<"mul24.lo.s32", B32, B32, B32, int_nvvm_mul24_i>;
def INT_NVVM_MUL24_UI : F_MATH_2<"mul24.lo.u32", B32, B32, B32, int_nvvm_mul24_ui>;
-def INT_NVVM_MUL_RN_SAT_F16 : F_MATH_2<"mul.rn.sat.f16", B16, B16, B16, int_nvvm_mul_rn_sat_f16>;
-def INT_NVVM_MUL_RN_FTZ_SAT_F16 : F_MATH_2<"mul.rn.ftz.sat.f16", B16, B16, B16, int_nvvm_mul_rn_ftz_sat_f16>;
-def INT_NVVM_MUL_RN_SAT_F16X2 : F_MATH_2<"mul.rn.sat.f16x2", B32, B32, B32, int_nvvm_mul_rn_sat_v2f16>;
-def INT_NVVM_MUL_RN_FTZ_SAT_F16X2 : F_MATH_2<"mul.rn.ftz.sat.f16x2", B32, B32, B32, int_nvvm_mul_rn_ftz_sat_v2f16>;
+// f16/f16x2
+foreach t = [F16RT, F16X2RT] in
+ foreach ftz = ["", "ftz"] in
+ foreach sat = ["", "sat"] in
+ def INT_NVVM_MUL_RN_ # !toupper(StrJoin<"_", [ftz, sat, t.PtxType]>.ret) :
+ F_MATH_2_RND_TY<StrJoin<".", ["mul.rn", ftz, sat, t.PtxType]>.ret, t,
+ !cast<Intrinsic>(
+ StrJoin<"_", ["int_nvvm_fmul", ftz, sat]>.ret),
+ fp_rnd_rn_imm>;
+
+// bf16/bf16x2
+foreach t = [BF16RT, BF16X2RT] in
+ def INT_NVVM_MUL_RN_ # !toupper(t.PtxType) :
+ F_MATH_2_RND_TY<"mul.rn." # t.PtxType, t, int_nvvm_fmul, fp_rnd_rn_imm,
+ BF16ArithPreds>;
+
+// f32/f32x2/f64
+foreach ftz = ["", "ftz"] in {
+ foreach sat = ["", "sat"] in
+ def StrJoin<"_", ["INT_NVVM_MUL", !toupper(ftz), !toupper(sat), "F"]>.ret :
+ F_MATH_2_RNDOP_TY<StrJoin<".", ["mul.${rnd}", ftz, sat, "f32"]>.ret,
+ F32RT,
+ !cast<Intrinsic>(
+ StrJoin<"_", ["int_nvvm_fmul", ftz, sat]>.ret)>;
+
+ def StrJoin<"_", ["INT_NVVM_MUL", !toupper(ftz), "F32X2"]>.ret :
+ F_MATH_2_RNDOP_TY<StrJoin<".", ["mul.${rnd}", ftz, "f32x2"]>.ret, F32X2RT,
+ !cast<Intrinsic>(
+ StrJoin<"_", ["int_nvvm_fmul", ftz]>.ret),
+ [hasF32x2Instructions]>;
+}
+
+def INT_NVVM_MUL_D :
+ F_MATH_2_RNDOP_TY<"mul.${rnd}.f64", F64RT, int_nvvm_fmul>;
//
// Div
@@ -2360,8 +2378,6 @@ def : Pat<(int_nvvm_cos_approx_f f32:$a), (COS_APPROX_f32 $a, NoFTZ)>;
// Fma
//
-defvar FPRoundingModes = ["rn", "rz", "rm", "rp"];
-
class FMA_TUPLE<string V, Intrinsic I, NVPTXRegClass RC,
list<Predicate> Preds = []> {
string Variant = V;
@@ -2569,16 +2585,6 @@ let Predicates = [doRsqrtOpt] in {
// Add
//
-defvar BF16ArithPreds = [hasBF16Math, PTX78, SM90];
-
-class FPRndModeImm<string mode> : TImmLeaf<i32,
- "return Imm == static_cast<int>(RoundingMode::" # mode # ");">;
-
-def fp_rnd_rn_imm : FPRndModeImm<"NearestTiesToEven">;
-def fp_rnd_rz_imm : FPRndModeImm<"TowardZero">;
-def fp_rnd_rm_imm : FPRndModeImm<"TowardNegative">;
-def fp_rnd_rp_imm : FPRndModeImm<"TowardPositive">;
-
// f16/f16x2
foreach t = [F16RT, F16X2RT] in
foreach ftz = ["", "ftz"] in
diff --git a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
index 05afba0452bc80..5f8f4899fc12a9 100644
--- a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
+++ b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
@@ -782,6 +782,30 @@ define void @nvvm_add(float %a, double %b, half %c, <2 x half> %d) {
ret void
}
+define void @nvvm_mul(float %a, double %b, half %c, <2 x half> %d) {
+; CHECK: call float @llvm.nvvm.fmul.f32(float %a, float %a, /* rnd=rn */ i32 1)
+; CHECK: call float @llvm.nvvm.fmul.ftz.f32(float %a, float %a, /* rnd=rz */ i32 0)
+; CHECK: call float @llvm.nvvm.fmul.f32(float %a, float %a, /* rnd=rm */ i32 3)
+; CHECK: call float @llvm.nvvm.fmul.ftz.f32(float %a, float %a, /* rnd=rp */ i32 2)
+; CHECK: call double @llvm.nvvm.fmul.f64(double %b, double %b, /* rnd=rn */ i32 1)
+; CHECK: call double @llvm.nvvm.fmul.f64(double %b, double %b, /* rnd=rz */ i32 0)
+; CHECK: call half @llvm.nvvm.fmul.sat.f16(half %c, half %c, /* rnd=rn */ i32 1)
+; CHECK: call half @llvm.nvvm.fmul.ftz.sat.f16(half %c, half %c, /* rnd=rn */ i32 1)
+; CHECK: call <2 x half> @llvm.nvvm.fmul.sat.v2f16(<2 x half> %d, <2 x half> %d, /* rnd=rn */ i32 1)
+; CHECK: call <2 x half> @llvm.nvvm.fmul.ftz.sat.v2f16(<2 x half> %d, <2 x half> %d, /* rnd=rn */ i32 1)
+ %r1 = call float @llvm.nvvm.mul.rn.f(float %a, float %a)
+ %r2 = call float @llvm.nvvm.mul.rz.ftz.f(float %a, float %a)
+ %r3 = call float @llvm.nvvm.mul.rm.f(float %a, float %a)
+ %r4 = call float @llvm.nvvm.mul.rp.ftz.f(float %a, float %a)
+ %r5 = call double @llvm.nvvm.mul.rn.d(double %b, double %b)
+ %r6 = call double @llvm.nvvm.mul.rz.d(double %b, double %b)
+ %r7 = call half @llvm.nvvm.mul.rn.sat.f16(half %c, half %c)
+ %r8 = call half @llvm.nvvm.mul.rn.ftz.sat.f16(half %c, half %c)
+ %r9 = call <2 x half> @llvm.nvvm.mul.rn.sat.v2f16(<2 x half> %d, <2 x half> %d)
+ %r10 = call <2 x half> @llvm.nvvm.mul.rn.ftz.sat.v2f16(<2 x half> %d, <2 x half> %d)
+ ret void
+}
+
declare void @llvm.nvvm.mbarrier.init(ptr, i32)
declare void @llvm.nvvm.mbarrier.init.shared(ptr addrspace(3), i32)
diff --git a/llvm/test/CodeGen/NVPTX/bf16-mul.ll b/llvm/test/CodeGen/NVPTX/bf16-mul.ll
new file mode 100644
index 00000000000000..3e877ebe29be4a
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/bf16-mul.ll
@@ -0,0 +1,33 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 | FileCheck %s
+; RUN: %if ptxas-sm_90 && ptxas-isa-7.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 | %ptxas-verify -arch=sm_90 %}
+
+define bfloat @mul_rn_bf16(bfloat %a, bfloat %b) {
+; CHECK-LABEL: mul_rn_bf16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [mul_rn_bf16_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [mul_rn_bf16_param_1];
+; CHECK-NEXT: mul.rn.bf16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;
+; CHECK-NEXT: ret;
+ %1 = call bfloat @llvm.nvvm.fmul.bf16(bfloat %a, bfloat %b, i32 1)
+ ret bfloat %1
+}
+
+define <2 x bfloat> @mul_rn_bf16x2(<2 x bfloat> %a, <2 x bfloat> %b) {
+; CHECK-LABEL: mul_rn_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [mul_rn_bf16x2_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [mul_rn_bf16x2_param_1];
+; CHECK-NEXT: mul.rn.bf16x2 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %1 = call <2 x bfloat> @llvm.nvvm.fmul.v2bf16(<2 x bfloat> %a, <2 x bfloat> %b, i32 1)
+ ret <2 x bfloat> %1
+}
diff --git a/llvm/test/CodeGen/NVPTX/f16-mul-sat.ll b/llvm/test/CodeGen/NVPTX/f16-mul-sat.ll
deleted file mode 100644
index 4bcc018f290d7a..00000000000000
--- a/llvm/test/CodeGen/NVPTX/f16-mul-sat.ll
+++ /dev/null
@@ -1,63 +0,0 @@
-; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
-; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_53 -mattr=+ptx42 | FileCheck %s
-; RUN: %if ptxas-isa-4.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_53 -mattr=+ptx42 | %ptxas-verify%}
-
-define half @mul_rn_sat_f16(half %a, half %b) {
-; CHECK-LABEL: mul_rn_sat_f16(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<4>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [mul_rn_sat_f16_param_0];
-; CHECK-NEXT: ld.param.b16 %rs2, [mul_rn_sat_f16_param_1];
-; CHECK-NEXT: mul.rn.sat.f16 %rs3, %rs1, %rs2;
-; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;
-; CHECK-NEXT: ret;
- %1 = call half @llvm.nvvm.mul.rn.sat.f16(half %a, half %b)
- ret half %1
-}
-
-define <2 x half> @mul_rn_sat_f16x2(<2 x half> %a, <2 x half> %b) {
-; CHECK-LABEL: mul_rn_sat_f16x2(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [mul_rn_sat_f16x2_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [mul_rn_sat_f16x2_param_1];
-; CHECK-NEXT: mul.rn.sat.f16x2 %r3, %r1, %r2;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
-; CHECK-NEXT: ret;
- %1 = call <2 x half> @llvm.nvvm.mul.rn.sat.v2f16(<2 x half> %a, <2 x half> %b)
- ret <2 x half> %1
-}
-
-define half @mul_rn_ftz_sat_f16(half %a, half %b) {
-; CHECK-LABEL: mul_rn_ftz_sat_f16(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<4>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [mul_rn_ftz_sat_f16_param_0];
-; CHECK-NEXT: ld.param.b16 %rs2, [mul_rn_ftz_sat_f16_param_1];
-; CHECK-NEXT: mul.rn.ftz.sat.f16 %rs3, %rs1, %rs2;
-; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;
-; CHECK-NEXT: ret;
- %1 = call half @llvm.nvvm.mul.rn.ftz.sat.f16(half %a, half %b)
- ret half %1
-}
-
-define <2 x half> @mul_rn_ftz_sat_f16x2(<2 x half> %a, <2 x half> %b) {
-; CHECK-LABEL: mul_rn_ftz_sat_f16x2(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [mul_rn_ftz_sat_f16x2_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [mul_rn_ftz_sat_f16x2_param_1];
-; CHECK-NEXT: mul.rn.ftz.sat.f16x2 %r3, %r1, %r2;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
-; CHECK-NEXT: ret;
- %1 = call <2 x half> @llvm.nvvm.mul.rn.ftz.sat.v2f16(<2 x half> %a, <2 x half> %b)
- ret <2 x half> %1
-}
diff --git a/llvm/test/CodeGen/NVPTX/f16-mul.ll b/llvm/test/CodeGen/NVPTX/f16-mul.ll
new file mode 100644
index 00000000000000..cbc8cadfce866e
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/f16-mul.ll
@@ -0,0 +1,123 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_53 -mattr=+ptx42 | FileCheck %s
+; RUN: %if ptxas-isa-4.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_53 -mattr=+ptx42 | %ptxas-verify%}
+
+define half @mul_rn_f16(half %a, half %b) {
+; CHECK-LABEL: mul_rn_f16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [mul_rn_f16_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [mul_rn_f16_param_1];
+; CHECK-NEXT: mul.rn.f16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;
+; CHECK-NEXT: ret;
+ %1 = call half @llvm.nvvm.fmul.f16(half %a, half %b, i32 1)
+ ret half %1
+}
+
+define <2 x half> @mul_rn_f16x2(<2 x half> %a, <2 x half> %b) {
+; CHECK-LABEL: mul_rn_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [mul_rn_f16x2_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [mul_rn_f16x2_param_1];
+; CHECK-NEXT: mul.rn.f16x2 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %1 = call <2 x half> @llvm.nvvm.fmul.v2f16(<2 x half> %a, <2 x half> %b, i32 1)
+ ret <2 x half> %1
+}
+
+define half @mul_rn_ftz_f16(half %a, half %b) {
+; CHECK-LABEL: mul_rn_ftz_f16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [mul_rn_ftz_f16_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [mul_rn_ftz_f16_param_1];
+; CHECK-NEXT: mul.rn.ftz.f16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;
+; CHECK-NEXT: ret;
+ %1 = call half @llvm.nvvm.fmul.ftz.f16(half %a, half %b, i32 1)
+ ret half %1
+}
+
+define <2 x half> @mul_rn_ftz_f16x2(<2 x half> %a, <2 x half> %b) {
+; CHECK-LABEL: mul_rn_ftz_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [mul_rn_ftz_f16x2_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [mul_rn_ftz_f16x2_param_1];
+; CHECK-NEXT: mul.rn.ftz.f16x2 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %1 = call <2 x half> @llvm.nvvm.fmul.ftz.v2f16(<2 x half> %a, <2 x half> %b, i32 1)
+ ret <2 x half> %1
+}
+
+define half @mul_rn_sat_f16(half %a, half %b) {
+; CHECK-LABEL: mul_rn_sat_f16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [mul_rn_sat_f16_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [mul_rn_sat_f16_param_1];
+; CHECK-NEXT: mul.rn.sat.f16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;
+; CHECK-NEXT: ret;
+ %1 = call half @llvm.nvvm.fmul.sat.f16(half %a, half %b, i32 1)
+ ret half %1
+}
+
+define <2 x half> @mul_rn_sat_f16x2(<2 x half> %a, <2 x half> %b) {
+; CHECK-LABEL: mul_rn_sat_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [mul_rn_sat_f16x2_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [mul_rn_sat_f16x2_param_1];
+; CHECK-NEXT: mul.rn.sat.f16x2 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %1 = call <2 x half> @llvm.nvvm.fmul.sat.v2f16(<2 x half> %a, <2 x half> %b, i32 1)
+ ret <2 x half> %1
+}
+
+define half @mul_rn_ftz_sat_f16(half %a, half %b) {
+; CHECK-LABEL: mul_rn_ftz_sat_f16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [mul_rn_ftz_sat_f16_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [mul_rn_ftz_sat_f16_param_1];
+; CHECK-NEXT: mul.rn.ftz.sat.f16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;
+; CHECK-NEXT: ret;
+ %1 = call half @llvm.nvvm.fmul.ftz.sat.f16(half %a, half %b, i32 1)
+ ret half %1
+}
+
+define <2 x half> @mul_rn_ftz_sat_f16x2(<2 x half> %a, <2 x half> %b) {
+; CHECK-LABEL: mul_rn_ftz_sat_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [mul_rn_ftz_sat_f16x2_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [mul_rn_ftz_sat_f16x2_param_1];
+; CHECK-NEXT: mul.rn.ftz.sat.f16x2 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %1 = call <2 x half> @llvm.nvvm.fmul.ftz.sat.v2f16(<2 x half> %a, <2 x half> %b, i32 1)
+ ret <2 x half> %1
+}
diff --git a/llvm/test/CodeGen/NVPTX/fp-arith-sat.ll b/llvm/test/CodeGen/NVPTX/fp-arith-sat.ll
index 0b93da7af17d02..8d1a00ae87eddc 100644
--- a/llvm/test/CodeGen/NVPTX/fp-arith-sat.ll
+++ b/llvm/test/CodeGen/NVPTX/fp-arith-sat.ll
@@ -80,6 +80,39 @@ define float @sub_sat_f32(float %a, float %b) {
ret float %r8
}
+define float @mul_sat_f32(float %a, float %b) {
+; CHECK-LABEL: mul_sat_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<11>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [mul_sat_f32_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [mul_sat_f32_param_1];
+; CHECK-NEXT: mul.rn.sat.f32 %r3, %r1, %r2;
+; CHECK-NEXT: mul.rn.ftz.sat.f32 %r4, %r1, %r3;
+; CHECK-NEXT: mul.rz.sat.f32 %r5, %r1, %r4;
+; CHECK-NEXT: mul.rz.ftz.sat.f32 %r6, %r1, %r5;
+; CHECK-NEXT: mul.rm.sat.f32 %r7, %r1, %r6;
+; CHECK-NEXT: mul.rm.ftz.sat.f32 %r8, %r1, %r7;
+; CHECK-NEXT: mul.rp.sat.f32 %r9, %r1, %r8;
+; CHECK-NEXT: mul.rp.ftz.sat.f32 %r10, %r1, %r9;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r10;
+; CHECK-NEXT: ret;
+ %r1 = call float @llvm.nvvm.fmul.sat.f32(float %a, float %b, i32 1)
+ %r2 = call float @llvm.nvvm.fmul.ftz.sat.f32(float %a, float %r1, i32 1)
+
+ %r3 = call float @llvm.nvvm.fmul.sat.f32(float %a, float %r2, i32 0)
+ %r4 = call float @llvm.nvvm.fmul.ftz.sat.f32(float %a, float %r3, i32 0)
+
+ %r5 = call float @llvm.nvvm.fmul.sat.f32(float %a, float %r4, i32 3)
+ %r6 = call float @llvm.nvvm.fmul.ftz.sat.f32(float %a, float %r5, i32 3)
+
+ %r7 = call float @llvm.nvvm.fmul.sat.f32(float %a, float %r6, i32 2)
+ %r8 = call float @llvm.nvvm.fmul.ftz.sat.f32(float %a, float %r7, i32 2)
+
+ ret float %r8
+}
+
define float @fma_sat_f32(float %a, float %b, float %c) {
; CHECK-LABEL: fma_sat_f32(
; CHECK: {
diff --git a/llvm/test/CodeGen/NVPTX/fp-mul-f32x2.ll b/llvm/test/CodeGen/NVPTX/fp-mul-f32x2.ll
new file mode 100644
index 00000000000000..9ad13c86ad6af2
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/fp-mul-f32x2.ll
@@ -0,0 +1,125 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mcpu=sm_100 -mattr=+ptx88 -march=nvptx64 | FileCheck %s
+; RUN: %if ptxas-sm_100 && ptxas-isa-8.8 %{ llc < %s -mcpu=sm_100 -mattr=+ptx88 -march=nvptx64 | %ptxas-verify -arch=sm_100 %}
+
+target triple = "nvptx64-nvidia-cuda"
+
+define <2 x float> @mul_rn(<2 x float> %a, <2 x float> %b) {
+; CHECK-LABEL: mul_rn(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [mul_rn_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [mul_rn_param_1];
+; CHECK-NEXT: mul.rn.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = call <2 x float> @llvm.nvvm.fmul.v2f32(<2 x float> %a, <2 x float> %b, i32 1)
+ ret <2 x float> %r
+}
+
+define <2 x float> @mul_rz(<2 x float> %a, <2 x float> %b) {
+; CHECK-LABEL: mul_rz(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [mul_rz_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [mul_rz_param_1];
+; CHECK-NEXT: mul.rz.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = call <2 x float> @llvm.nvvm.fmul.v2f32(<2 x float> %a, <2 x float> %b, i32 0)
+ ret <2 x float> %r
+}
+
+define <2 x float> @mul_rm(<2 x float> %a, <2 x float> %b) {
+; CHECK-LABEL: mul_rm(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [mul_rm_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [mul_rm_param_1];
+; CHECK-NEXT: mul.rm.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = call <2 x float> @llvm.nvvm.fmul.v2f32(<2 x float> %a, <2 x float> %b, i32 3)
+ ret <2 x float> %r
+}
+
+define <2 x float> @mul_rp(<2 x float> %a, <2 x float> %b) {
+; CHECK-LABEL: mul_rp(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [mul_rp_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [mul_rp_param_1];
+; CHECK-NEXT: mul.rp.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = call <2 x float> @llvm.nvvm.fmul.v2f32(<2 x float> %a, <2 x float> %b, i32 2)
+ ret <2 x float> %r
+}
+
+define <2 x float> @mul_rn_ftz(<2 x float> %a, <2 x float> %b) {
+; CHECK-LABEL: mul_rn_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [mul_rn_ftz_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [mul_rn_ftz_param_1];
+; CHECK-NEXT: mul.rn.ftz.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = call <2 x float> @llvm.nvvm.fmul.ftz.v2f32(<2 x float> %a, <2 x float> %b, i32 1)
+ ret <2 x float> %r
+}
+
+define <2 x float> @mul_rz_ftz(<2 x float> %a, <2 x float> %b) {
+; CHECK-LABEL: mul_rz_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [mul_rz_ftz_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [mul_rz_ftz_param_1];
+; CHECK-NEXT: mul.rz.ftz.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = call <2 x float> @llvm.nvvm.fmul.ftz.v2f32(<2 x float> %a, <2 x float> %b, i32 0)
+ ret <2 x float> %r
+}
+
+define <2 x float> @mul_rm_ftz(<2 x float> %a, <2 x float> %b) {
+; CHECK-LABEL: mul_rm_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [mul_rm_ftz_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [mul_rm_ftz_param_1];
+; CHECK-NEXT: mul.rm.ftz.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = call <2 x float> @llvm.nvvm.fmul.ftz.v2f32(<2 x float> %a, <2 x float> %b, i32 3)
+ ret <2 x float> %r
+}
+
+define <2 x float> @mul_rp_ftz(<2 x float> %a, <2 x float> %b) {
+; CHECK-LABEL: mul_rp_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b64 %rd1, [mul_rp_ftz_param_0];
+; CHECK-NEXT: ld.param::func.b64 %rd2, [mul_rp_ftz_param_1];
+; CHECK-NEXT: mul.rp.ftz.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param::func.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = call <2 x float> @llvm.nvvm.fmul.ftz.v2f32(<2 x float> %a, <2 x float> %b, i32 2)
+ ret <2 x float> %r
+}
diff --git a/llvm/test/CodeGen/NVPTX/fp-mul-invalid.ll b/llvm/test/CodeGen/NVPTX/fp-mul-invalid.ll
new file mode 100644
index 00000000000000..2ae57c156e28b3
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/fp-mul-invalid.ll
@@ -0,0 +1,47 @@
+; RUN: not llc < %s -mcpu=sm_100 -mattr=+ptx88 -march=nvptx64 2>&1 | FileCheck %s
+; RUN: not llc < %s -mcpu=sm_90 -mattr=+ptx78 -march=nvptx64 2>&1 | FileCheck %s --check-prefix=NOF32X2
+; RUN: not llc < %s -mcpu=sm_80 -mattr=+ptx78 -march=nvptx64 2>&1 | FileCheck %s --check-prefix=NOBF16
+
+target triple = "nvptx64-nvidia-cuda"
+
+; CHECK: error: {{.*}}llvm.nvvm.fmul.sat with rounding mode rn and operand type v2f32 is not supported
+define <2 x float> @sat_f32x2(<2 x float> %a, <2 x float> %b) {
+ %r = call <2 x float> @llvm.nvvm.fmul.sat.v2f32(<2 x float> %a, <2 x float> %b, i32 1)
+ ret <2 x float> %r
+}
+
+; NOF32X2: error: {{.*}}llvm.nvvm.fmul with rounding mode rn and operand type v2f32 is not supported
+define <2 x float> @unsupported_f32x2(<2 x float> %a, <2 x float> %b) {
+ %r = call <2 x float> @llvm.nvvm.fmul.v2f32(<2 x float> %a, <2 x float> %b, i32 1)
+ ret <2 x float> %r
+}
+
+; CHECK: error: {{.*}}llvm.nvvm.fmul.ftz with rounding mode rn and operand type f64 is not supported
+define double @ftz_f64(double %a, double %b) {
+ %r = call double @llvm.nvvm.fmul.ftz.f64(double %a, double %b, i32 1)
+ ret double %r
+}
+
+; CHECK: error: {{.*}}llvm.nvvm.fmul.sat with rounding mode rz and operand type f16 is not supported
+define half @rz_f16(half %a, half %b) {
+ %r = call half @llvm.nvvm.fmul.sat.f16(half %a, half %b, i32 0)
+ ret half %r
+}
+
+; CHECK: error: {{.*}}llvm.nvvm.fmul.ftz with rounding mode rn and operand type bf16 is not supported
+define bfloat @ftz_bf16(bfloat %a, bfloat %b) {
+ %r = call bfloat @llvm.nvvm.fmul.ftz.bf16(bfloat %a, bfloat %b, i32 1)
+ ret bfloat %r
+}
+
+; NOBF16: error: {{.*}}llvm.nvvm.fmul with rounding mode rn and operand type bf16 is not supported
+define bfloat @unsupported_bf16(bfloat %a, bfloat %b) {
+ %r = call bfloat @llvm.nvvm.fmul.bf16(bfloat %a, bfloat %b, i32 1)
+ ret bfloat %r
+}
+
+; CHECK: error: {{.*}}llvm.nvvm.fmul with rounding mode rn and operand type v4f32 is not supported
+define <4 x float> @v4f32(<4 x float> %a, <4 x float> %b) {
+ %r = call <4 x float> @llvm.nvvm.fmul.v4f32(<4 x float> %a, <4 x float> %b, i32 1)
+ ret <4 x float> %r
+}
diff --git a/llvm/test/CodeGen/NVPTX/fp-mul.ll b/llvm/test/CodeGen/NVPTX/fp-mul.ll
new file mode 100644
index 00000000000000..262d7307b7f22f
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/fp-mul.ll
@@ -0,0 +1,58 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s
+; RUN: %if ptxas-sm_20 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify -arch=sm_20 %}
+
+define float @mul_f32(float %a, float %b) {
+; CHECK-LABEL: mul_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<11>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [mul_f32_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [mul_f32_param_1];
+; CHECK-NEXT: mul.rn.f32 %r3, %r1, %r2;
+; CHECK-NEXT: mul.rn.ftz.f32 %r4, %r1, %r3;
+; CHECK-NEXT: mul.rz.f32 %r5, %r1, %r4;
+; CHECK-NEXT: mul.rz.ftz.f32 %r6, %r1, %r5;
+; CHECK-NEXT: mul.rm.f32 %r7, %r1, %r6;
+; CHECK-NEXT: mul.rm.ftz.f32 %r8, %r1, %r7;
+; CHECK-NEXT: mul.rp.f32 %r9, %r1, %r8;
+; CHECK-NEXT: mul.rp.ftz.f32 %r10, %r1, %r9;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r10;
+; CHECK-NEXT: ret;
+ %r1 = call float @llvm.nvvm.fmul.f32(float %a, float %b, i32 1)
+ %r2 = call float @llvm.nvvm.fmul.ftz.f32(float %a, float %r1, i32 1)
+
+ %r3 = call float @llvm.nvvm.fmul.f32(float %a, float %r2, i32 0)
+ %r4 = call float @llvm.nvvm.fmul.ftz.f32(float %a, float %r3, i32 0)
+
+ %r5 = call float @llvm.nvvm.fmul.f32(float %a, float %r4, i32 3)
+ %r6 = call float @llvm.nvvm.fmul.ftz.f32(float %a, float %r5, i32 3)
+
+ %r7 = call float @llvm.nvvm.fmul.f32(float %a, float %r6, i32 2)
+ %r8 = call float @llvm.nvvm.fmul.ftz.f32(float %a, float %r7, i32 2)
+
+ ret float %r8
+}
+
+define double @mul_f64(double %a, double %b) {
+; CHECK-LABEL: mul_f64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [mul_f64_param_0];
+; CHECK-NEXT: ld.param.b64 %rd2, [mul_f64_param_1];
+; CHECK-NEXT: mul.rn.f64 %rd3, %rd1, %rd2;
+; CHECK-NEXT: mul.rz.f64 %rd4, %rd1, %rd3;
+; CHECK-NEXT: mul.rm.f64 %rd5, %rd1, %rd4;
+; CHECK-NEXT: mul.rp.f64 %rd6, %rd1, %rd5;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd6;
+; CHECK-NEXT: ret;
+ %r1 = call double @llvm.nvvm.fmul.f64(double %a, double %b, i32 1)
+ %r2 = call double @llvm.nvvm.fmul.f64(double %a, double %r1, i32 0)
+ %r3 = call double @llvm.nvvm.fmul.f64(double %a, double %r2, i32 3)
+ %r4 = call double @llvm.nvvm.fmul.f64(double %a, double %r3, i32 2)
+
+ ret double %r4
+}
diff --git a/llvm/test/Transforms/InstCombine/NVPTX/nvvm-intrins.ll b/llvm/test/Transforms/InstCombine/NVPTX/nvvm-intrins.ll
index 8e48f6f3c7c1cc..ff3727e3bbb563 100644
--- a/llvm/test/Transforms/InstCombine/NVPTX/nvvm-intrins.ll
+++ b/llvm/test/Transforms/InstCombine/NVPTX/nvvm-intrins.ll
@@ -318,20 +318,20 @@ define float @test_add_rn_f_ftz(float %a, float %b) #0 {
; CHECK-LABEL: @test_mul_rn_d
define double @test_mul_rn_d(double %a, double %b) #0 {
-; CHECK: call double @llvm.nvvm.mul.rn.d
- %ret = call double @llvm.nvvm.mul.rn.d(double %a, double %b)
+; CHECK: call double @llvm.nvvm.fmul.f64
+ %ret = call double @llvm.nvvm.fmul.f64(double %a, double %b, /* rnd=rn */ i32 1)
ret double %ret
}
; CHECK-LABEL: @test_mul_rn_f
define float @test_mul_rn_f(float %a, float %b) #0 {
-; CHECK: call float @llvm.nvvm.mul.rn.f
- %ret = call float @llvm.nvvm.mul.rn.f(float %a, float %b)
+; CHECK: call float @llvm.nvvm.fmul.f32
+ %ret = call float @llvm.nvvm.fmul.f32(float %a, float %b, /* rnd=rn */ i32 1)
ret float %ret
}
; CHECK-LABEL: @test_mul_rn_f_ftz
define float @test_mul_rn_f_ftz(float %a, float %b) #0 {
-; CHECK: call float @llvm.nvvm.mul.rn.ftz.f(float %a, float %b)
- %ret = call float @llvm.nvvm.mul.rn.ftz.f(float %a, float %b)
+; CHECK: call float @llvm.nvvm.fmul.ftz.f32(float %a, float %b, /* rnd=rn */ i32 1)
+ %ret = call float @llvm.nvvm.fmul.ftz.f32(float %a, float %b, /* rnd=rn */ i32 1)
ret float %ret
}
@@ -513,9 +513,9 @@ declare float @llvm.nvvm.i2f.rn(i32)
declare double @llvm.nvvm.ll2d.rn(i64)
declare float @llvm.nvvm.ll2f.rn(i64)
declare double @llvm.nvvm.lohi.i2d(i32, i32)
-declare double @llvm.nvvm.mul.rn.d(double, double)
-declare float @llvm.nvvm.mul.rn.f(float, float)
-declare float @llvm.nvvm.mul.rn.ftz.f(float, float)
+declare double @llvm.nvvm.fmul.f64(double, double, i32 immarg)
+declare float @llvm.nvvm.fmul.f32(float, float, i32 immarg)
+declare float @llvm.nvvm.fmul.ftz.f32(float, float, i32 immarg)
declare double @llvm.nvvm.rcp.rm.d(double)
declare double @llvm.nvvm.rcp.rn.d(double)
declare float @llvm.nvvm.rcp.rn.f(float)
diff --git a/llvm/test/Transforms/InstSimplify/const-fold-nvvm-mul.ll b/llvm/test/Transforms/InstSimplify/const-fold-nvvm-mul.ll
index e643c02524d1de..37af66b50dc432 100644
--- a/llvm/test/Transforms/InstSimplify/const-fold-nvvm-mul.ll
+++ b/llvm/test/Transforms/InstSimplify/const-fold-nvvm-mul.ll
@@ -13,7 +13,7 @@ define double @test_1_25_times_2_rm_d() {
; CHECK-LABEL: define double @test_1_25_times_2_rm_d() {
; CHECK-NEXT: ret double 2.500000e+00
;
- %res = call double @llvm.nvvm.mul.rm.d(double 1.25, double 2.0)
+ %res = call double @llvm.nvvm.fmul.f64(double 1.25, double 2.0, /* rnd=rm */ i32 3)
ret double %res
}
@@ -21,7 +21,7 @@ define double @test_1_25_times_2_rn_d() {
; CHECK-LABEL: define double @test_1_25_times_2_rn_d() {
; CHECK-NEXT: ret double 2.500000e+00
;
- %res = call double @llvm.nvvm.mul.rn.d(double 1.25, double 2.0)
+ %res = call double @llvm.nvvm.fmul.f64(double 1.25, double 2.0, /* rnd=rn */ i32 1)
ret double %res
}
@@ -29,7 +29,7 @@ define double @test_1_25_times_2_rp_d() {
; CHECK-LABEL: define double @test_1_25_times_2_rp_d() {
; CHECK-NEXT: ret double 2.500000e+00
;
- %res = call double @llvm.nvvm.mul.rp.d(double 1.25, double 2.0)
+ %res = call double @llvm.nvvm.fmul.f64(double 1.25, double 2.0, /* rnd=rp */ i32 2)
ret double %res
}
@@ -37,7 +37,7 @@ define double @test_1_25_times_2_rz_d() {
; CHECK-LABEL: define double @test_1_25_times_2_rz_d() {
; CHECK-NEXT: ret double 2.500000e+00
;
- %res = call double @llvm.nvvm.mul.rz.d(double 1.25, double 2.0)
+ %res = call double @llvm.nvvm.fmul.f64(double 1.25, double 2.0, /* rnd=rz */ i32 0)
ret double %res
}
@@ -45,7 +45,7 @@ define float @test_1_25_times_2_rm_f() {
; CHECK-LABEL: define float @test_1_25_times_2_rm_f() {
; CHECK-NEXT: ret float 2.500000e+00
;
- %res = call float @llvm.nvvm.mul.rm.f(float 1.25, float 2.0)
+ %res = call float @llvm.nvvm.fmul.f32(float 1.25, float 2.0, /* rnd=rm */ i32 3)
ret float %res
}
@@ -53,7 +53,7 @@ define float @test_1_25_times_2_rn_f() {
; CHECK-LABEL: define float @test_1_25_times_2_rn_f() {
; CHECK-NEXT: ret float 2.500000e+00
;
- %res = call float @llvm.nvvm.mul.rn.f(float 1.25, float 2.0)
+ %res = call float @llvm.nvvm.fmul.f32(float 1.25, float 2.0, /* rnd=rn */ i32 1)
ret float %res
}
@@ -61,7 +61,7 @@ define float @test_1_25_times_2_rp_f() {
; CHECK-LABEL: define float @test_1_25_times_2_rp_f() {
; CHECK-NEXT: ret float 2.500000e+00
;
- %res = call float @llvm.nvvm.mul.rp.f(float 1.25, float 2.0)
+ %res = call float @llvm.nvvm.fmul.f32(float 1.25, float 2.0, /* rnd=rp */ i32 2)
ret float %res
}
@@ -69,7 +69,7 @@ define float @test_1_25_times_2_rz_f() {
; CHECK-LABEL: define float @test_1_25_times_2_rz_f() {
; CHECK-NEXT: ret float 2.500000e+00
;
- %res = call float @llvm.nvvm.mul.rz.f(float 1.25, float 2.0)
+ %res = call float @llvm.nvvm.fmul.f32(float 1.25, float 2.0, /* rnd=rz */ i32 0)
ret float %res
}
@@ -77,7 +77,7 @@ define float @test_1_25_times_2_rm_ftz_f() {
; CHECK-LABEL: define float @test_1_25_times_2_rm_ftz_f() {
; CHECK-NEXT: ret float 2.500000e+00
;
- %res = call float @llvm.nvvm.mul.rm.ftz.f(float 1.25, float 2.0)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 1.25, float 2.0, /* rnd=rm */ i32 3)
ret float %res
}
@@ -85,7 +85,7 @@ define float @test_1_25_times_2_rn_ftz_f() {
; CHECK-LABEL: define float @test_1_25_times_2_rn_ftz_f() {
; CHECK-NEXT: ret float 2.500000e+00
;
- %res = call float @llvm.nvvm.mul.rn.ftz.f(float 1.25, float 2.0)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 1.25, float 2.0, /* rnd=rn */ i32 1)
ret float %res
}
@@ -93,7 +93,7 @@ define float @test_1_25_times_2_rp_ftz_f() {
; CHECK-LABEL: define float @test_1_25_times_2_rp_ftz_f() {
; CHECK-NEXT: ret float 2.500000e+00
;
- %res = call float @llvm.nvvm.mul.rp.ftz.f(float 1.25, float 2.0)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 1.25, float 2.0, /* rnd=rp */ i32 2)
ret float %res
}
@@ -101,10 +101,74 @@ define float @test_1_25_times_2_rz_ftz_f() {
; CHECK-LABEL: define float @test_1_25_times_2_rz_ftz_f() {
; CHECK-NEXT: ret float 2.500000e+00
;
- %res = call float @llvm.nvvm.mul.rz.ftz.f(float 1.25, float 2.0)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 1.25, float 2.0, /* rnd=rz */ i32 0)
ret float %res
}
+define half @test_1_25_times_2_rm_f16() {
+; CHECK-LABEL: define half @test_1_25_times_2_rm_f16() {
+; CHECK-NEXT: ret half 2.500000e+00
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 1.25, half 2.0, /* rnd=rm */ i32 3)
+ ret half %res
+}
+
+define half @test_1_25_times_2_rn_f16() {
+; CHECK-LABEL: define half @test_1_25_times_2_rn_f16() {
+; CHECK-NEXT: ret half 2.500000e+00
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 1.25, half 2.0, /* rnd=rn */ i32 1)
+ ret half %res
+}
+
+define half @test_1_25_times_2_rp_f16() {
+; CHECK-LABEL: define half @test_1_25_times_2_rp_f16() {
+; CHECK-NEXT: ret half 2.500000e+00
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 1.25, half 2.0, /* rnd=rp */ i32 2)
+ ret half %res
+}
+
+define half @test_1_25_times_2_rz_f16() {
+; CHECK-LABEL: define half @test_1_25_times_2_rz_f16() {
+; CHECK-NEXT: ret half 2.500000e+00
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 1.25, half 2.0, /* rnd=rz */ i32 0)
+ ret half %res
+}
+
+define bfloat @test_1_25_times_2_rm_bf16() {
+; CHECK-LABEL: define bfloat @test_1_25_times_2_rm_bf16() {
+; CHECK-NEXT: ret bfloat 2.500000e+00
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 1.25, bfloat 2.0, /* rnd=rm */ i32 3)
+ ret bfloat %res
+}
+
+define bfloat @test_1_25_times_2_rn_bf16() {
+; CHECK-LABEL: define bfloat @test_1_25_times_2_rn_bf16() {
+; CHECK-NEXT: ret bfloat 2.500000e+00
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 1.25, bfloat 2.0, /* rnd=rn */ i32 1)
+ ret bfloat %res
+}
+
+define bfloat @test_1_25_times_2_rp_bf16() {
+; CHECK-LABEL: define bfloat @test_1_25_times_2_rp_bf16() {
+; CHECK-NEXT: ret bfloat 2.500000e+00
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 1.25, bfloat 2.0, /* rnd=rp */ i32 2)
+ ret bfloat %res
+}
+
+define bfloat @test_1_25_times_2_rz_bf16() {
+; CHECK-LABEL: define bfloat @test_1_25_times_2_rz_bf16() {
+; CHECK-NEXT: ret bfloat 2.500000e+00
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 1.25, bfloat 2.0, /* rnd=rz */ i32 0)
+ ret bfloat %res
+}
+
;###############################################################
;# Mul(1.0, Subnormal) #
;###############################################################
@@ -118,7 +182,7 @@ define double @test_1_times_subnorm_rm_d() {
; CHECK-LABEL: define double @test_1_times_subnorm_rm_d() {
; CHECK-NEXT: ret double 4.940660e-324
;
- %res = call double @llvm.nvvm.mul.rm.d(double 1.0, double 0x0000000000000001)
+ %res = call double @llvm.nvvm.fmul.f64(double 1.0, double 0x0000000000000001, /* rnd=rm */ i32 3)
ret double %res
}
@@ -126,7 +190,7 @@ define double @test_1_times_subnorm_rn_d() {
; CHECK-LABEL: define double @test_1_times_subnorm_rn_d() {
; CHECK-NEXT: ret double 4.940660e-324
;
- %res = call double @llvm.nvvm.mul.rn.d(double 1.0, double 0x0000000000000001)
+ %res = call double @llvm.nvvm.fmul.f64(double 1.0, double 0x0000000000000001, /* rnd=rn */ i32 1)
ret double %res
}
@@ -134,7 +198,7 @@ define double @test_1_times_subnorm_rp_d() {
; CHECK-LABEL: define double @test_1_times_subnorm_rp_d() {
; CHECK-NEXT: ret double 4.940660e-324
;
- %res = call double @llvm.nvvm.mul.rp.d(double 1.0, double 0x0000000000000001)
+ %res = call double @llvm.nvvm.fmul.f64(double 1.0, double 0x0000000000000001, /* rnd=rp */ i32 2)
ret double %res
}
@@ -142,7 +206,7 @@ define double @test_1_times_subnorm_rz_d() {
; CHECK-LABEL: define double @test_1_times_subnorm_rz_d() {
; CHECK-NEXT: ret double 4.940660e-324
;
- %res = call double @llvm.nvvm.mul.rz.d(double 1.0, double 0x0000000000000001)
+ %res = call double @llvm.nvvm.fmul.f64(double 1.0, double 0x0000000000000001, /* rnd=rz */ i32 0)
ret double %res
}
@@ -150,7 +214,7 @@ define float @test_1_times_subnorm_rm_f() {
; CHECK-LABEL: define float @test_1_times_subnorm_rm_f() {
; CHECK-NEXT: ret float 1.401300e-45
;
- %res = call float @llvm.nvvm.mul.rm.f(float 1.0, float 0x36A0000000000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 1.0, float 0x36A0000000000000, /* rnd=rm */ i32 3)
ret float %res
}
@@ -158,7 +222,7 @@ define float @test_1_times_subnorm_rn_f() {
; CHECK-LABEL: define float @test_1_times_subnorm_rn_f() {
; CHECK-NEXT: ret float 1.401300e-45
;
- %res = call float @llvm.nvvm.mul.rn.f(float 1.0, float 0x36A0000000000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 1.0, float 0x36A0000000000000, /* rnd=rn */ i32 1)
ret float %res
}
@@ -166,7 +230,7 @@ define float @test_1_times_subnorm_rp_f() {
; CHECK-LABEL: define float @test_1_times_subnorm_rp_f() {
; CHECK-NEXT: ret float 1.401300e-45
;
- %res = call float @llvm.nvvm.mul.rp.f(float 1.0, float 0x36A0000000000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 1.0, float 0x36A0000000000000, /* rnd=rp */ i32 2)
ret float %res
}
@@ -174,7 +238,7 @@ define float @test_1_times_subnorm_rz_f() {
; CHECK-LABEL: define float @test_1_times_subnorm_rz_f() {
; CHECK-NEXT: ret float 1.401300e-45
;
- %res = call float @llvm.nvvm.mul.rz.f(float 1.0, float 0x36A0000000000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 1.0, float 0x36A0000000000000, /* rnd=rz */ i32 0)
ret float %res
}
@@ -182,7 +246,7 @@ define float @test_1_times_subnorm_rm_ftz_f() {
; CHECK-LABEL: define float @test_1_times_subnorm_rm_ftz_f() {
; CHECK-NEXT: ret float 0.000000e+00
;
- %res = call float @llvm.nvvm.mul.rm.ftz.f(float 1.0, float 0x36A0000000000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 1.0, float 0x36A0000000000000, /* rnd=rm */ i32 3)
ret float %res
}
@@ -190,7 +254,7 @@ define float @test_1_times_subnorm_rn_ftz_f() {
; CHECK-LABEL: define float @test_1_times_subnorm_rn_ftz_f() {
; CHECK-NEXT: ret float 0.000000e+00
;
- %res = call float @llvm.nvvm.mul.rn.ftz.f(float 1.0, float 0x36A0000000000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 1.0, float 0x36A0000000000000, /* rnd=rn */ i32 1)
ret float %res
}
@@ -198,7 +262,7 @@ define float @test_1_times_subnorm_rp_ftz_f() {
; CHECK-LABEL: define float @test_1_times_subnorm_rp_ftz_f() {
; CHECK-NEXT: ret float 0.000000e+00
;
- %res = call float @llvm.nvvm.mul.rp.ftz.f(float 1.0, float 0x36A0000000000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 1.0, float 0x36A0000000000000, /* rnd=rp */ i32 2)
ret float %res
}
@@ -206,10 +270,106 @@ define float @test_1_times_subnorm_rz_ftz_f() {
; CHECK-LABEL: define float @test_1_times_subnorm_rz_ftz_f() {
; CHECK-NEXT: ret float 0.000000e+00
;
- %res = call float @llvm.nvvm.mul.rz.ftz.f(float 1.0, float 0x36A0000000000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 1.0, float 0x36A0000000000000, /* rnd=rz */ i32 0)
ret float %res
}
+define half @test_1_times_subnorm_rm_f16() {
+; CHECK-LABEL: define half @test_1_times_subnorm_rm_f16() {
+; CHECK-NEXT: ret half 5.960460e-08
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 1.0, half 0xH0001, /* rnd=rm */ i32 3)
+ ret half %res
+}
+
+define half @test_1_times_subnorm_rn_f16() {
+; CHECK-LABEL: define half @test_1_times_subnorm_rn_f16() {
+; CHECK-NEXT: ret half 5.960460e-08
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 1.0, half 0xH0001, /* rnd=rn */ i32 1)
+ ret half %res
+}
+
+define half @test_1_times_subnorm_rp_f16() {
+; CHECK-LABEL: define half @test_1_times_subnorm_rp_f16() {
+; CHECK-NEXT: ret half 5.960460e-08
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 1.0, half 0xH0001, /* rnd=rp */ i32 2)
+ ret half %res
+}
+
+define half @test_1_times_subnorm_rz_f16() {
+; CHECK-LABEL: define half @test_1_times_subnorm_rz_f16() {
+; CHECK-NEXT: ret half 5.960460e-08
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 1.0, half 0xH0001, /* rnd=rz */ i32 0)
+ ret half %res
+}
+
+define half @test_1_times_subnorm_rm_ftz_f16() {
+; CHECK-LABEL: define half @test_1_times_subnorm_rm_ftz_f16() {
+; CHECK-NEXT: ret half 0.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.ftz.f16(half 1.0, half 0xH0001, /* rnd=rm */ i32 3)
+ ret half %res
+}
+
+define half @test_1_times_subnorm_rn_ftz_f16() {
+; CHECK-LABEL: define half @test_1_times_subnorm_rn_ftz_f16() {
+; CHECK-NEXT: ret half 0.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.ftz.f16(half 1.0, half 0xH0001, /* rnd=rn */ i32 1)
+ ret half %res
+}
+
+define half @test_1_times_subnorm_rp_ftz_f16() {
+; CHECK-LABEL: define half @test_1_times_subnorm_rp_ftz_f16() {
+; CHECK-NEXT: ret half 0.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.ftz.f16(half 1.0, half 0xH0001, /* rnd=rp */ i32 2)
+ ret half %res
+}
+
+define half @test_1_times_subnorm_rz_ftz_f16() {
+; CHECK-LABEL: define half @test_1_times_subnorm_rz_ftz_f16() {
+; CHECK-NEXT: ret half 0.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.ftz.f16(half 1.0, half 0xH0001, /* rnd=rz */ i32 0)
+ ret half %res
+}
+
+define bfloat @test_1_times_subnorm_rm_bf16() {
+; CHECK-LABEL: define bfloat @test_1_times_subnorm_rm_bf16() {
+; CHECK-NEXT: ret bfloat 9.183550e-41
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 1.0, bfloat 0xR0001, /* rnd=rm */ i32 3)
+ ret bfloat %res
+}
+
+define bfloat @test_1_times_subnorm_rn_bf16() {
+; CHECK-LABEL: define bfloat @test_1_times_subnorm_rn_bf16() {
+; CHECK-NEXT: ret bfloat 9.183550e-41
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 1.0, bfloat 0xR0001, /* rnd=rn */ i32 1)
+ ret bfloat %res
+}
+
+define bfloat @test_1_times_subnorm_rp_bf16() {
+; CHECK-LABEL: define bfloat @test_1_times_subnorm_rp_bf16() {
+; CHECK-NEXT: ret bfloat 9.183550e-41
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 1.0, bfloat 0xR0001, /* rnd=rp */ i32 2)
+ ret bfloat %res
+}
+
+define bfloat @test_1_times_subnorm_rz_bf16() {
+; CHECK-LABEL: define bfloat @test_1_times_subnorm_rz_bf16() {
+; CHECK-NEXT: ret bfloat 9.183550e-41
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 1.0, bfloat 0xR0001, /* rnd=rz */ i32 0)
+ ret bfloat %res
+}
+
;###############################################################
;# Mul(1.0, -Subnormal) #
;###############################################################
@@ -223,7 +383,7 @@ define double @test_1_times_neg_subnorm_rm_d() {
; CHECK-LABEL: define double @test_1_times_neg_subnorm_rm_d() {
; CHECK-NEXT: ret double -4.940660e-324
;
- %res = call double @llvm.nvvm.mul.rm.d(double 1.0, double 0x8000000000000001)
+ %res = call double @llvm.nvvm.fmul.f64(double 1.0, double 0x8000000000000001, /* rnd=rm */ i32 3)
ret double %res
}
@@ -231,7 +391,7 @@ define double @test_1_times_neg_subnorm_rn_d() {
; CHECK-LABEL: define double @test_1_times_neg_subnorm_rn_d() {
; CHECK-NEXT: ret double -4.940660e-324
;
- %res = call double @llvm.nvvm.mul.rn.d(double 1.0, double 0x8000000000000001)
+ %res = call double @llvm.nvvm.fmul.f64(double 1.0, double 0x8000000000000001, /* rnd=rn */ i32 1)
ret double %res
}
@@ -239,7 +399,7 @@ define double @test_1_times_neg_subnorm_rp_d() {
; CHECK-LABEL: define double @test_1_times_neg_subnorm_rp_d() {
; CHECK-NEXT: ret double -4.940660e-324
;
- %res = call double @llvm.nvvm.mul.rp.d(double 1.0, double 0x8000000000000001)
+ %res = call double @llvm.nvvm.fmul.f64(double 1.0, double 0x8000000000000001, /* rnd=rp */ i32 2)
ret double %res
}
@@ -247,7 +407,7 @@ define double @test_1_times_neg_subnorm_rz_d() {
; CHECK-LABEL: define double @test_1_times_neg_subnorm_rz_d() {
; CHECK-NEXT: ret double -4.940660e-324
;
- %res = call double @llvm.nvvm.mul.rz.d(double 1.0, double 0x8000000000000001)
+ %res = call double @llvm.nvvm.fmul.f64(double 1.0, double 0x8000000000000001, /* rnd=rz */ i32 0)
ret double %res
}
@@ -255,7 +415,7 @@ define float @test_1_times_neg_subnorm_rm_f() {
; CHECK-LABEL: define float @test_1_times_neg_subnorm_rm_f() {
; CHECK-NEXT: ret float -1.401300e-45
;
- %res = call float @llvm.nvvm.mul.rm.f(float 1.0, float 0xB6A0000000000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 1.0, float 0xB6A0000000000000, /* rnd=rm */ i32 3)
ret float %res
}
@@ -263,7 +423,7 @@ define float @test_1_times_neg_subnorm_rn_f() {
; CHECK-LABEL: define float @test_1_times_neg_subnorm_rn_f() {
; CHECK-NEXT: ret float -1.401300e-45
;
- %res = call float @llvm.nvvm.mul.rn.f(float 1.0, float 0xB6A0000000000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 1.0, float 0xB6A0000000000000, /* rnd=rn */ i32 1)
ret float %res
}
@@ -271,7 +431,7 @@ define float @test_1_times_neg_subnorm_rp_f() {
; CHECK-LABEL: define float @test_1_times_neg_subnorm_rp_f() {
; CHECK-NEXT: ret float -1.401300e-45
;
- %res = call float @llvm.nvvm.mul.rp.f(float 1.0, float 0xB6A0000000000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 1.0, float 0xB6A0000000000000, /* rnd=rp */ i32 2)
ret float %res
}
@@ -279,7 +439,7 @@ define float @test_1_times_neg_subnorm_rz_f() {
; CHECK-LABEL: define float @test_1_times_neg_subnorm_rz_f() {
; CHECK-NEXT: ret float -1.401300e-45
;
- %res = call float @llvm.nvvm.mul.rz.f(float 1.0, float 0xB6A0000000000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 1.0, float 0xB6A0000000000000, /* rnd=rz */ i32 0)
ret float %res
}
@@ -287,7 +447,7 @@ define float @test_1_times_neg_subnorm_rm_ftz_f() {
; CHECK-LABEL: define float @test_1_times_neg_subnorm_rm_ftz_f() {
; CHECK-NEXT: ret float -0.000000e+00
;
- %res = call float @llvm.nvvm.mul.rm.ftz.f(float 1.0, float 0xB6A0000000000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 1.0, float 0xB6A0000000000000, /* rnd=rm */ i32 3)
ret float %res
}
@@ -295,7 +455,7 @@ define float @test_1_times_neg_subnorm_rn_ftz_f() {
; CHECK-LABEL: define float @test_1_times_neg_subnorm_rn_ftz_f() {
; CHECK-NEXT: ret float -0.000000e+00
;
- %res = call float @llvm.nvvm.mul.rn.ftz.f(float 1.0, float 0xB6A0000000000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 1.0, float 0xB6A0000000000000, /* rnd=rn */ i32 1)
ret float %res
}
@@ -303,7 +463,7 @@ define float @test_1_times_neg_subnorm_rp_ftz_f() {
; CHECK-LABEL: define float @test_1_times_neg_subnorm_rp_ftz_f() {
; CHECK-NEXT: ret float -0.000000e+00
;
- %res = call float @llvm.nvvm.mul.rp.ftz.f(float 1.0, float 0xB6A0000000000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 1.0, float 0xB6A0000000000000, /* rnd=rp */ i32 2)
ret float %res
}
@@ -311,10 +471,106 @@ define float @test_1_times_neg_subnorm_rz_ftz_f() {
; CHECK-LABEL: define float @test_1_times_neg_subnorm_rz_ftz_f() {
; CHECK-NEXT: ret float -0.000000e+00
;
- %res = call float @llvm.nvvm.mul.rz.ftz.f(float 1.0, float 0xB6A0000000000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 1.0, float 0xB6A0000000000000, /* rnd=rz */ i32 0)
ret float %res
}
+define half @test_1_times_neg_subnorm_rm_f16() {
+; CHECK-LABEL: define half @test_1_times_neg_subnorm_rm_f16() {
+; CHECK-NEXT: ret half -5.960460e-08
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 1.0, half 0xH8001, /* rnd=rm */ i32 3)
+ ret half %res
+}
+
+define half @test_1_times_neg_subnorm_rn_f16() {
+; CHECK-LABEL: define half @test_1_times_neg_subnorm_rn_f16() {
+; CHECK-NEXT: ret half -5.960460e-08
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 1.0, half 0xH8001, /* rnd=rn */ i32 1)
+ ret half %res
+}
+
+define half @test_1_times_neg_subnorm_rp_f16() {
+; CHECK-LABEL: define half @test_1_times_neg_subnorm_rp_f16() {
+; CHECK-NEXT: ret half -5.960460e-08
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 1.0, half 0xH8001, /* rnd=rp */ i32 2)
+ ret half %res
+}
+
+define half @test_1_times_neg_subnorm_rz_f16() {
+; CHECK-LABEL: define half @test_1_times_neg_subnorm_rz_f16() {
+; CHECK-NEXT: ret half -5.960460e-08
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 1.0, half 0xH8001, /* rnd=rz */ i32 0)
+ ret half %res
+}
+
+define half @test_1_times_neg_subnorm_rm_ftz_f16() {
+; CHECK-LABEL: define half @test_1_times_neg_subnorm_rm_ftz_f16() {
+; CHECK-NEXT: ret half -0.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.ftz.f16(half 1.0, half 0xH8001, /* rnd=rm */ i32 3)
+ ret half %res
+}
+
+define half @test_1_times_neg_subnorm_rn_ftz_f16() {
+; CHECK-LABEL: define half @test_1_times_neg_subnorm_rn_ftz_f16() {
+; CHECK-NEXT: ret half -0.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.ftz.f16(half 1.0, half 0xH8001, /* rnd=rn */ i32 1)
+ ret half %res
+}
+
+define half @test_1_times_neg_subnorm_rp_ftz_f16() {
+; CHECK-LABEL: define half @test_1_times_neg_subnorm_rp_ftz_f16() {
+; CHECK-NEXT: ret half -0.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.ftz.f16(half 1.0, half 0xH8001, /* rnd=rp */ i32 2)
+ ret half %res
+}
+
+define half @test_1_times_neg_subnorm_rz_ftz_f16() {
+; CHECK-LABEL: define half @test_1_times_neg_subnorm_rz_ftz_f16() {
+; CHECK-NEXT: ret half -0.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.ftz.f16(half 1.0, half 0xH8001, /* rnd=rz */ i32 0)
+ ret half %res
+}
+
+define bfloat @test_1_times_neg_subnorm_rm_bf16() {
+; CHECK-LABEL: define bfloat @test_1_times_neg_subnorm_rm_bf16() {
+; CHECK-NEXT: ret bfloat -9.183550e-41
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 1.0, bfloat 0xR8001, /* rnd=rm */ i32 3)
+ ret bfloat %res
+}
+
+define bfloat @test_1_times_neg_subnorm_rn_bf16() {
+; CHECK-LABEL: define bfloat @test_1_times_neg_subnorm_rn_bf16() {
+; CHECK-NEXT: ret bfloat -9.183550e-41
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 1.0, bfloat 0xR8001, /* rnd=rn */ i32 1)
+ ret bfloat %res
+}
+
+define bfloat @test_1_times_neg_subnorm_rp_bf16() {
+; CHECK-LABEL: define bfloat @test_1_times_neg_subnorm_rp_bf16() {
+; CHECK-NEXT: ret bfloat -9.183550e-41
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 1.0, bfloat 0xR8001, /* rnd=rp */ i32 2)
+ ret bfloat %res
+}
+
+define bfloat @test_1_times_neg_subnorm_rz_bf16() {
+; CHECK-LABEL: define bfloat @test_1_times_neg_subnorm_rz_bf16() {
+; CHECK-NEXT: ret bfloat -9.183550e-41
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 1.0, bfloat 0xR8001, /* rnd=rz */ i32 0)
+ ret bfloat %res
+}
+
;###############################################################
;# Mul(Normal, Normal) -> Subnormal #
;###############################################################
@@ -327,7 +583,7 @@ define double @test_normal_times_normal_to_subnorm_rm_d() {
; CHECK-LABEL: define double @test_normal_times_normal_to_subnorm_rm_d() {
; CHECK-NEXT: ret double f0x3800000000000000
;
- %res = call double @llvm.nvvm.mul.rm.d(double 0x3810000000000000, double 0.5)
+ %res = call double @llvm.nvvm.fmul.f64(double 0x3810000000000000, double 0.5, /* rnd=rm */ i32 3)
ret double %res
}
@@ -335,7 +591,7 @@ define double @test_normal_times_normal_to_subnorm_rn_d() {
; CHECK-LABEL: define double @test_normal_times_normal_to_subnorm_rn_d() {
; CHECK-NEXT: ret double f0x3800000000000000
;
- %res = call double @llvm.nvvm.mul.rn.d(double 0x3810000000000000, double 0.5)
+ %res = call double @llvm.nvvm.fmul.f64(double 0x3810000000000000, double 0.5, /* rnd=rn */ i32 1)
ret double %res
}
@@ -343,7 +599,7 @@ define double @test_normal_times_normal_to_subnorm_rp_d() {
; CHECK-LABEL: define double @test_normal_times_normal_to_subnorm_rp_d() {
; CHECK-NEXT: ret double f0x3800000000000000
;
- %res = call double @llvm.nvvm.mul.rp.d(double 0x3810000000000000, double 0.5)
+ %res = call double @llvm.nvvm.fmul.f64(double 0x3810000000000000, double 0.5, /* rnd=rp */ i32 2)
ret double %res
}
@@ -351,7 +607,7 @@ define double @test_normal_times_normal_to_subnorm_rz_d() {
; CHECK-LABEL: define double @test_normal_times_normal_to_subnorm_rz_d() {
; CHECK-NEXT: ret double f0x3800000000000000
;
- %res = call double @llvm.nvvm.mul.rz.d(double 0x3810000000000000, double 0.5)
+ %res = call double @llvm.nvvm.fmul.f64(double 0x3810000000000000, double 0.5, /* rnd=rz */ i32 0)
ret double %res
}
@@ -359,7 +615,7 @@ define float @test_normal_times_normal_to_subnorm_rm_f() {
; CHECK-LABEL: define float @test_normal_times_normal_to_subnorm_rm_f() {
; CHECK-NEXT: ret float f0x00400000
;
- %res = call float @llvm.nvvm.mul.rm.f(float 0x3810000000000000, float 0.5)
+ %res = call float @llvm.nvvm.fmul.f32(float 0x3810000000000000, float 0.5, /* rnd=rm */ i32 3)
ret float %res
}
@@ -367,7 +623,7 @@ define float @test_normal_times_normal_to_subnorm_rn_f() {
; CHECK-LABEL: define float @test_normal_times_normal_to_subnorm_rn_f() {
; CHECK-NEXT: ret float f0x00400000
;
- %res = call float @llvm.nvvm.mul.rn.f(float 0x3810000000000000, float 0.5)
+ %res = call float @llvm.nvvm.fmul.f32(float 0x3810000000000000, float 0.5, /* rnd=rn */ i32 1)
ret float %res
}
@@ -375,7 +631,7 @@ define float @test_normal_times_normal_to_subnorm_rp_f() {
; CHECK-LABEL: define float @test_normal_times_normal_to_subnorm_rp_f() {
; CHECK-NEXT: ret float f0x00400000
;
- %res = call float @llvm.nvvm.mul.rp.f(float 0x3810000000000000, float 0.5)
+ %res = call float @llvm.nvvm.fmul.f32(float 0x3810000000000000, float 0.5, /* rnd=rp */ i32 2)
ret float %res
}
@@ -383,7 +639,7 @@ define float @test_normal_times_normal_to_subnorm_rz_f() {
; CHECK-LABEL: define float @test_normal_times_normal_to_subnorm_rz_f() {
; CHECK-NEXT: ret float f0x00400000
;
- %res = call float @llvm.nvvm.mul.rz.f(float 0x3810000000000000, float 0.5)
+ %res = call float @llvm.nvvm.fmul.f32(float 0x3810000000000000, float 0.5, /* rnd=rz */ i32 0)
ret float %res
}
@@ -391,7 +647,7 @@ define float @test_normal_times_normal_to_subnorm_rm_ftz_f() {
; CHECK-LABEL: define float @test_normal_times_normal_to_subnorm_rm_ftz_f() {
; CHECK-NEXT: ret float 0.000000e+00
;
- %res = call float @llvm.nvvm.mul.rm.ftz.f(float 0x3810000000000000, float 0.5)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0x3810000000000000, float 0.5, /* rnd=rm */ i32 3)
ret float %res
}
@@ -399,7 +655,7 @@ define float @test_normal_times_normal_to_subnorm_rn_ftz_f() {
; CHECK-LABEL: define float @test_normal_times_normal_to_subnorm_rn_ftz_f() {
; CHECK-NEXT: ret float 0.000000e+00
;
- %res = call float @llvm.nvvm.mul.rn.ftz.f(float 0x3810000000000000, float 0.5)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0x3810000000000000, float 0.5, /* rnd=rn */ i32 1)
ret float %res
}
@@ -407,7 +663,7 @@ define float @test_normal_times_normal_to_subnorm_rp_ftz_f() {
; CHECK-LABEL: define float @test_normal_times_normal_to_subnorm_rp_ftz_f() {
; CHECK-NEXT: ret float 0.000000e+00
;
- %res = call float @llvm.nvvm.mul.rp.ftz.f(float 0x3810000000000000, float 0.5)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0x3810000000000000, float 0.5, /* rnd=rp */ i32 2)
ret float %res
}
@@ -415,10 +671,106 @@ define float @test_normal_times_normal_to_subnorm_rz_ftz_f() {
; CHECK-LABEL: define float @test_normal_times_normal_to_subnorm_rz_ftz_f() {
; CHECK-NEXT: ret float 0.000000e+00
;
- %res = call float @llvm.nvvm.mul.rz.ftz.f(float 0x3810000000000000, float 0.5)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0x3810000000000000, float 0.5, /* rnd=rz */ i32 0)
ret float %res
}
+define half @test_normal_times_normal_to_subnorm_rm_f16() {
+; CHECK-LABEL: define half @test_normal_times_normal_to_subnorm_rm_f16() {
+; CHECK-NEXT: ret half 3.051760e-05
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 0xH0400, half 0.5, /* rnd=rm */ i32 3)
+ ret half %res
+}
+
+define half @test_normal_times_normal_to_subnorm_rn_f16() {
+; CHECK-LABEL: define half @test_normal_times_normal_to_subnorm_rn_f16() {
+; CHECK-NEXT: ret half 3.051760e-05
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 0xH0400, half 0.5, /* rnd=rn */ i32 1)
+ ret half %res
+}
+
+define half @test_normal_times_normal_to_subnorm_rp_f16() {
+; CHECK-LABEL: define half @test_normal_times_normal_to_subnorm_rp_f16() {
+; CHECK-NEXT: ret half 3.051760e-05
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 0xH0400, half 0.5, /* rnd=rp */ i32 2)
+ ret half %res
+}
+
+define half @test_normal_times_normal_to_subnorm_rz_f16() {
+; CHECK-LABEL: define half @test_normal_times_normal_to_subnorm_rz_f16() {
+; CHECK-NEXT: ret half 3.051760e-05
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 0xH0400, half 0.5, /* rnd=rz */ i32 0)
+ ret half %res
+}
+
+define half @test_normal_times_normal_to_subnorm_rm_ftz_f16() {
+; CHECK-LABEL: define half @test_normal_times_normal_to_subnorm_rm_ftz_f16() {
+; CHECK-NEXT: ret half 0.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.ftz.f16(half 0xH0400, half 0.5, /* rnd=rm */ i32 3)
+ ret half %res
+}
+
+define half @test_normal_times_normal_to_subnorm_rn_ftz_f16() {
+; CHECK-LABEL: define half @test_normal_times_normal_to_subnorm_rn_ftz_f16() {
+; CHECK-NEXT: ret half 0.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.ftz.f16(half 0xH0400, half 0.5, /* rnd=rn */ i32 1)
+ ret half %res
+}
+
+define half @test_normal_times_normal_to_subnorm_rp_ftz_f16() {
+; CHECK-LABEL: define half @test_normal_times_normal_to_subnorm_rp_ftz_f16() {
+; CHECK-NEXT: ret half 0.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.ftz.f16(half 0xH0400, half 0.5, /* rnd=rp */ i32 2)
+ ret half %res
+}
+
+define half @test_normal_times_normal_to_subnorm_rz_ftz_f16() {
+; CHECK-LABEL: define half @test_normal_times_normal_to_subnorm_rz_ftz_f16() {
+; CHECK-NEXT: ret half 0.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.ftz.f16(half 0xH0400, half 0.5, /* rnd=rz */ i32 0)
+ ret half %res
+}
+
+define bfloat @test_normal_times_normal_to_subnorm_rm_bf16() {
+; CHECK-LABEL: define bfloat @test_normal_times_normal_to_subnorm_rm_bf16() {
+; CHECK-NEXT: ret bfloat 5.877470e-39
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 0xR0080, bfloat 0.5, /* rnd=rm */ i32 3)
+ ret bfloat %res
+}
+
+define bfloat @test_normal_times_normal_to_subnorm_rn_bf16() {
+; CHECK-LABEL: define bfloat @test_normal_times_normal_to_subnorm_rn_bf16() {
+; CHECK-NEXT: ret bfloat 5.877470e-39
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 0xR0080, bfloat 0.5, /* rnd=rn */ i32 1)
+ ret bfloat %res
+}
+
+define bfloat @test_normal_times_normal_to_subnorm_rp_bf16() {
+; CHECK-LABEL: define bfloat @test_normal_times_normal_to_subnorm_rp_bf16() {
+; CHECK-NEXT: ret bfloat 5.877470e-39
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 0xR0080, bfloat 0.5, /* rnd=rp */ i32 2)
+ ret bfloat %res
+}
+
+define bfloat @test_normal_times_normal_to_subnorm_rz_bf16() {
+; CHECK-LABEL: define bfloat @test_normal_times_normal_to_subnorm_rz_bf16() {
+; CHECK-NEXT: ret bfloat 5.877470e-39
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 0xR0080, bfloat 0.5, /* rnd=rz */ i32 0)
+ ret bfloat %res
+}
+
;###############################################################
;# Mul(2.0, NaN) #
;###############################################################
@@ -427,112 +779,184 @@ define float @test_normal_times_normal_to_subnorm_rz_ftz_f() {
define double @test_2_times_nan_rm_d() {
; CHECK-LABEL: define double @test_2_times_nan_rm_d() {
-; CHECK-NEXT: [[RES:%.*]] = call double @llvm.nvvm.mul.rm.d(double 2.000000e+00, double +snan(0x4444400000000))
+; CHECK-NEXT: [[RES:%.*]] = call double @llvm.nvvm.fmul.f64(double 2.000000e+00, double +snan(0x4444400000000), /* rnd=rm */ i32 3)
; CHECK-NEXT: ret double [[RES]]
;
- %res = call double @llvm.nvvm.mul.rm.d(double 2.0, double 0x7FF4444400000000)
+ %res = call double @llvm.nvvm.fmul.f64(double 2.0, double 0x7FF4444400000000, /* rnd=rm */ i32 3)
ret double %res
}
define double @test_2_times_nan_rn_d() {
; CHECK-LABEL: define double @test_2_times_nan_rn_d() {
-; CHECK-NEXT: [[RES:%.*]] = call double @llvm.nvvm.mul.rn.d(double 2.000000e+00, double +snan(0x4444400000000))
+; CHECK-NEXT: [[RES:%.*]] = call double @llvm.nvvm.fmul.f64(double 2.000000e+00, double +snan(0x4444400000000), /* rnd=rn */ i32 1)
; CHECK-NEXT: ret double [[RES]]
;
- %res = call double @llvm.nvvm.mul.rn.d(double 2.0, double 0x7FF4444400000000)
+ %res = call double @llvm.nvvm.fmul.f64(double 2.0, double 0x7FF4444400000000, /* rnd=rn */ i32 1)
ret double %res
}
define double @test_2_times_nan_rp_d() {
; CHECK-LABEL: define double @test_2_times_nan_rp_d() {
-; CHECK-NEXT: [[RES:%.*]] = call double @llvm.nvvm.mul.rp.d(double 2.000000e+00, double +snan(0x4444400000000))
+; CHECK-NEXT: [[RES:%.*]] = call double @llvm.nvvm.fmul.f64(double 2.000000e+00, double +snan(0x4444400000000), /* rnd=rp */ i32 2)
; CHECK-NEXT: ret double [[RES]]
;
- %res = call double @llvm.nvvm.mul.rp.d(double 2.0, double 0x7FF4444400000000)
+ %res = call double @llvm.nvvm.fmul.f64(double 2.0, double 0x7FF4444400000000, /* rnd=rp */ i32 2)
ret double %res
}
define double @test_2_times_nan_rz_d() {
; CHECK-LABEL: define double @test_2_times_nan_rz_d() {
-; CHECK-NEXT: [[RES:%.*]] = call double @llvm.nvvm.mul.rz.d(double 2.000000e+00, double +snan(0x4444400000000))
+; CHECK-NEXT: [[RES:%.*]] = call double @llvm.nvvm.fmul.f64(double 2.000000e+00, double +snan(0x4444400000000), /* rnd=rz */ i32 0)
; CHECK-NEXT: ret double [[RES]]
;
- %res = call double @llvm.nvvm.mul.rz.d(double 2.0, double 0x7FF4444400000000)
+ %res = call double @llvm.nvvm.fmul.f64(double 2.0, double 0x7FF4444400000000, /* rnd=rz */ i32 0)
ret double %res
}
define float @test_2_times_nan_rm_f() {
; CHECK-LABEL: define float @test_2_times_nan_rm_f() {
-; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.mul.rm.f(float 2.000000e+00, float +nan(0x3A2220))
+; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.fmul.f32(float 2.000000e+00, float +nan(0x3A2220), /* rnd=rm */ i32 3)
; CHECK-NEXT: ret float [[RES]]
;
- %res = call float @llvm.nvvm.mul.rm.f(float 2.0, float 0x7FFF444400000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 2.0, float 0x7FFF444400000000, /* rnd=rm */ i32 3)
ret float %res
}
define float @test_2_times_nan_rn_f() {
; CHECK-LABEL: define float @test_2_times_nan_rn_f() {
-; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.mul.rn.f(float 2.000000e+00, float +nan(0x3A2220))
+; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.fmul.f32(float 2.000000e+00, float +nan(0x3A2220), /* rnd=rn */ i32 1)
; CHECK-NEXT: ret float [[RES]]
;
- %res = call float @llvm.nvvm.mul.rn.f(float 2.0, float 0x7FFF444400000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 2.0, float 0x7FFF444400000000, /* rnd=rn */ i32 1)
ret float %res
}
define float @test_2_times_nan_rp_f() {
; CHECK-LABEL: define float @test_2_times_nan_rp_f() {
-; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.mul.rp.f(float 2.000000e+00, float +nan(0x3A2220))
+; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.fmul.f32(float 2.000000e+00, float +nan(0x3A2220), /* rnd=rp */ i32 2)
; CHECK-NEXT: ret float [[RES]]
;
- %res = call float @llvm.nvvm.mul.rp.f(float 2.0, float 0x7FFF444400000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 2.0, float 0x7FFF444400000000, /* rnd=rp */ i32 2)
ret float %res
}
define float @test_2_times_nan_rz_f() {
; CHECK-LABEL: define float @test_2_times_nan_rz_f() {
-; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.mul.rz.f(float 2.000000e+00, float +nan(0x3A2220))
+; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.fmul.f32(float 2.000000e+00, float +nan(0x3A2220), /* rnd=rz */ i32 0)
; CHECK-NEXT: ret float [[RES]]
;
- %res = call float @llvm.nvvm.mul.rz.f(float 2.0, float 0x7FFF444400000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 2.0, float 0x7FFF444400000000, /* rnd=rz */ i32 0)
ret float %res
}
define float @test_2_times_nan_rm_ftz_f() {
; CHECK-LABEL: define float @test_2_times_nan_rm_ftz_f() {
-; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.mul.rm.ftz.f(float 2.000000e+00, float +nan(0x3A2220))
+; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.fmul.ftz.f32(float 2.000000e+00, float +nan(0x3A2220), /* rnd=rm */ i32 3)
; CHECK-NEXT: ret float [[RES]]
;
- %res = call float @llvm.nvvm.mul.rm.ftz.f(float 2.0, float 0x7FFF444400000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 2.0, float 0x7FFF444400000000, /* rnd=rm */ i32 3)
ret float %res
}
define float @test_2_times_nan_rn_ftz_f() {
; CHECK-LABEL: define float @test_2_times_nan_rn_ftz_f() {
-; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.mul.rn.ftz.f(float 2.000000e+00, float +nan(0x3A2220))
+; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.fmul.ftz.f32(float 2.000000e+00, float +nan(0x3A2220), /* rnd=rn */ i32 1)
; CHECK-NEXT: ret float [[RES]]
;
- %res = call float @llvm.nvvm.mul.rn.ftz.f(float 2.0, float 0x7FFF444400000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 2.0, float 0x7FFF444400000000, /* rnd=rn */ i32 1)
ret float %res
}
define float @test_2_times_nan_rp_ftz_f() {
; CHECK-LABEL: define float @test_2_times_nan_rp_ftz_f() {
-; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.mul.rp.ftz.f(float 2.000000e+00, float +nan(0x3A2220))
+; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.fmul.ftz.f32(float 2.000000e+00, float +nan(0x3A2220), /* rnd=rp */ i32 2)
; CHECK-NEXT: ret float [[RES]]
;
- %res = call float @llvm.nvvm.mul.rp.ftz.f(float 2.0, float 0x7FFF444400000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 2.0, float 0x7FFF444400000000, /* rnd=rp */ i32 2)
ret float %res
}
define float @test_2_times_nan_rz_ftz_f() {
; CHECK-LABEL: define float @test_2_times_nan_rz_ftz_f() {
-; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.mul.rz.ftz.f(float 2.000000e+00, float +nan(0x3A2220))
+; CHECK-NEXT: [[RES:%.*]] = call float @llvm.nvvm.fmul.ftz.f32(float 2.000000e+00, float +nan(0x3A2220), /* rnd=rz */ i32 0)
; CHECK-NEXT: ret float [[RES]]
;
- %res = call float @llvm.nvvm.mul.rz.ftz.f(float 2.0, float 0x7FFF444400000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 2.0, float 0x7FFF444400000000, /* rnd=rz */ i32 0)
ret float %res
}
+define half @test_2_times_nan_rm_f16() {
+; CHECK-LABEL: define half @test_2_times_nan_rm_f16() {
+; CHECK-NEXT: [[RES:%.*]] = call half @llvm.nvvm.fmul.f16(half 2.000000e+00, half +qnan, /* rnd=rm */ i32 3)
+; CHECK-NEXT: ret half [[RES]]
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 2.0, half 0xH7E00, /* rnd=rm */ i32 3)
+ ret half %res
+}
+
+define half @test_2_times_nan_rn_f16() {
+; CHECK-LABEL: define half @test_2_times_nan_rn_f16() {
+; CHECK-NEXT: [[RES:%.*]] = call half @llvm.nvvm.fmul.f16(half 2.000000e+00, half +qnan, /* rnd=rn */ i32 1)
+; CHECK-NEXT: ret half [[RES]]
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 2.0, half 0xH7E00, /* rnd=rn */ i32 1)
+ ret half %res
+}
+
+define half @test_2_times_nan_rp_f16() {
+; CHECK-LABEL: define half @test_2_times_nan_rp_f16() {
+; CHECK-NEXT: [[RES:%.*]] = call half @llvm.nvvm.fmul.f16(half 2.000000e+00, half +qnan, /* rnd=rp */ i32 2)
+; CHECK-NEXT: ret half [[RES]]
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 2.0, half 0xH7E00, /* rnd=rp */ i32 2)
+ ret half %res
+}
+
+define half @test_2_times_nan_rz_f16() {
+; CHECK-LABEL: define half @test_2_times_nan_rz_f16() {
+; CHECK-NEXT: [[RES:%.*]] = call half @llvm.nvvm.fmul.f16(half 2.000000e+00, half +qnan, /* rnd=rz */ i32 0)
+; CHECK-NEXT: ret half [[RES]]
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 2.0, half 0xH7E00, /* rnd=rz */ i32 0)
+ ret half %res
+}
+
+define bfloat @test_2_times_nan_rm_bf16() {
+; CHECK-LABEL: define bfloat @test_2_times_nan_rm_bf16() {
+; CHECK-NEXT: [[RES:%.*]] = call bfloat @llvm.nvvm.fmul.bf16(bfloat 2.000000e+00, bfloat +qnan, /* rnd=rm */ i32 3)
+; CHECK-NEXT: ret bfloat [[RES]]
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 2.0, bfloat 0xR7FC0, /* rnd=rm */ i32 3)
+ ret bfloat %res
+}
+
+define bfloat @test_2_times_nan_rn_bf16() {
+; CHECK-LABEL: define bfloat @test_2_times_nan_rn_bf16() {
+; CHECK-NEXT: [[RES:%.*]] = call bfloat @llvm.nvvm.fmul.bf16(bfloat 2.000000e+00, bfloat +qnan, /* rnd=rn */ i32 1)
+; CHECK-NEXT: ret bfloat [[RES]]
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 2.0, bfloat 0xR7FC0, /* rnd=rn */ i32 1)
+ ret bfloat %res
+}
+
+define bfloat @test_2_times_nan_rp_bf16() {
+; CHECK-LABEL: define bfloat @test_2_times_nan_rp_bf16() {
+; CHECK-NEXT: [[RES:%.*]] = call bfloat @llvm.nvvm.fmul.bf16(bfloat 2.000000e+00, bfloat +qnan, /* rnd=rp */ i32 2)
+; CHECK-NEXT: ret bfloat [[RES]]
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 2.0, bfloat 0xR7FC0, /* rnd=rp */ i32 2)
+ ret bfloat %res
+}
+
+define bfloat @test_2_times_nan_rz_bf16() {
+; CHECK-LABEL: define bfloat @test_2_times_nan_rz_bf16() {
+; CHECK-NEXT: [[RES:%.*]] = call bfloat @llvm.nvvm.fmul.bf16(bfloat 2.000000e+00, bfloat +qnan, /* rnd=rz */ i32 0)
+; CHECK-NEXT: ret bfloat [[RES]]
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 2.0, bfloat 0xR7FC0, /* rnd=rz */ i32 0)
+ ret bfloat %res
+}
+
;###############################################################
;# Mul(0.75, 4/3 + epsilon) #
;###############################################################
@@ -547,7 +971,7 @@ define float @test_mul_just_above_1_rm_f() {
; CHECK-LABEL: define float @test_mul_just_above_1_rm_f() {
; CHECK-NEXT: ret float 1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rm.f(float 0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 0.75, float 0x3FF5555560000000, /* rnd=rm */ i32 3)
ret float %res
}
@@ -555,7 +979,7 @@ define float @test_mul_just_above_1_rn_f() {
; CHECK-LABEL: define float @test_mul_just_above_1_rn_f() {
; CHECK-NEXT: ret float 1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rn.f(float 0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 0.75, float 0x3FF5555560000000, /* rnd=rn */ i32 1)
ret float %res
}
@@ -563,7 +987,7 @@ define float @test_mul_just_above_1_rp_f() {
; CHECK-LABEL: define float @test_mul_just_above_1_rp_f() {
; CHECK-NEXT: ret float f0x3F800001
;
- %res = call float @llvm.nvvm.mul.rp.f(float 0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 0.75, float 0x3FF5555560000000, /* rnd=rp */ i32 2)
ret float %res
}
@@ -571,7 +995,7 @@ define float @test_mul_just_above_1_rz_f() {
; CHECK-LABEL: define float @test_mul_just_above_1_rz_f() {
; CHECK-NEXT: ret float 1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rz.f(float 0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 0.75, float 0x3FF5555560000000, /* rnd=rz */ i32 0)
ret float %res
}
@@ -579,7 +1003,7 @@ define float @test_mul_just_above_1_rm_ftz_f() {
; CHECK-LABEL: define float @test_mul_just_above_1_rm_ftz_f() {
; CHECK-NEXT: ret float 1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rm.ftz.f(float 0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0.75, float 0x3FF5555560000000, /* rnd=rm */ i32 3)
ret float %res
}
@@ -587,7 +1011,7 @@ define float @test_mul_just_above_1_rn_ftz_f() {
; CHECK-LABEL: define float @test_mul_just_above_1_rn_ftz_f() {
; CHECK-NEXT: ret float 1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rn.ftz.f(float 0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0.75, float 0x3FF5555560000000, /* rnd=rn */ i32 1)
ret float %res
}
@@ -595,7 +1019,7 @@ define float @test_mul_just_above_1_rp_ftz_f() {
; CHECK-LABEL: define float @test_mul_just_above_1_rp_ftz_f() {
; CHECK-NEXT: ret float f0x3F800001
;
- %res = call float @llvm.nvvm.mul.rp.ftz.f(float 0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0.75, float 0x3FF5555560000000, /* rnd=rp */ i32 2)
ret float %res
}
@@ -603,7 +1027,7 @@ define float @test_mul_just_above_1_rz_ftz_f() {
; CHECK-LABEL: define float @test_mul_just_above_1_rz_ftz_f() {
; CHECK-NEXT: ret float 1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rz.ftz.f(float 0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0.75, float 0x3FF5555560000000, /* rnd=rz */ i32 0)
ret float %res
}
@@ -621,7 +1045,7 @@ define double @test_mul_just_above_1_rm_d() {
; CHECK-LABEL: define double @test_mul_just_above_1_rm_d() {
; CHECK-NEXT: ret double 1.000000e+00
;
- %res = call double @llvm.nvvm.mul.rm.d(double 0.75, double 0x3FF5555555555556)
+ %res = call double @llvm.nvvm.fmul.f64(double 0.75, double 0x3FF5555555555556, /* rnd=rm */ i32 3)
ret double %res
}
@@ -629,7 +1053,7 @@ define double @test_mul_just_above_1_rn_d() {
; CHECK-LABEL: define double @test_mul_just_above_1_rn_d() {
; CHECK-NEXT: ret double 1.000000e+00
;
- %res = call double @llvm.nvvm.mul.rn.d(double 0.75, double 0x3FF5555555555556)
+ %res = call double @llvm.nvvm.fmul.f64(double 0.75, double 0x3FF5555555555556, /* rnd=rn */ i32 1)
ret double %res
}
@@ -637,7 +1061,7 @@ define double @test_mul_just_above_1_rp_d() {
; CHECK-LABEL: define double @test_mul_just_above_1_rp_d() {
; CHECK-NEXT: ret double f0x3FF0000000000001
;
- %res = call double @llvm.nvvm.mul.rp.d(double 0.75, double 0x3FF5555555555556)
+ %res = call double @llvm.nvvm.fmul.f64(double 0.75, double 0x3FF5555555555556, /* rnd=rp */ i32 2)
ret double %res
}
@@ -645,7 +1069,7 @@ define double @test_mul_just_above_1_rz_d() {
; CHECK-LABEL: define double @test_mul_just_above_1_rz_d() {
; CHECK-NEXT: ret double 1.000000e+00
;
- %res = call double @llvm.nvvm.mul.rz.d(double 0.75, double 0x3FF5555555555556)
+ %res = call double @llvm.nvvm.fmul.f64(double 0.75, double 0x3FF5555555555556, /* rnd=rz */ i32 0)
ret double %res
}
@@ -663,7 +1087,7 @@ define float @test_mul_just_below_negative_1_rm_f() {
; CHECK-LABEL: define float @test_mul_just_below_negative_1_rm_f() {
; CHECK-NEXT: ret float f0xBF800001
;
- %res = call float @llvm.nvvm.mul.rm.f(float -0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.f32(float -0.75, float 0x3FF5555560000000, /* rnd=rm */ i32 3)
ret float %res
}
@@ -671,7 +1095,7 @@ define float @test_mul_just_below_negative_1_rn_f() {
; CHECK-LABEL: define float @test_mul_just_below_negative_1_rn_f() {
; CHECK-NEXT: ret float -1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rn.f(float -0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.f32(float -0.75, float 0x3FF5555560000000, /* rnd=rn */ i32 1)
ret float %res
}
@@ -679,7 +1103,7 @@ define float @test_mul_just_below_negative_1_rp_f() {
; CHECK-LABEL: define float @test_mul_just_below_negative_1_rp_f() {
; CHECK-NEXT: ret float -1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rp.f(float -0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.f32(float -0.75, float 0x3FF5555560000000, /* rnd=rp */ i32 2)
ret float %res
}
@@ -687,7 +1111,7 @@ define float @test_mul_just_below_negative_1_rz_f() {
; CHECK-LABEL: define float @test_mul_just_below_negative_1_rz_f() {
; CHECK-NEXT: ret float -1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rz.f(float -0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.f32(float -0.75, float 0x3FF5555560000000, /* rnd=rz */ i32 0)
ret float %res
}
@@ -695,7 +1119,7 @@ define float @test_mul_just_below_negative_1_rm_ftz_f() {
; CHECK-LABEL: define float @test_mul_just_below_negative_1_rm_ftz_f() {
; CHECK-NEXT: ret float f0xBF800001
;
- %res = call float @llvm.nvvm.mul.rm.ftz.f(float -0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float -0.75, float 0x3FF5555560000000, /* rnd=rm */ i32 3)
ret float %res
}
@@ -703,7 +1127,7 @@ define float @test_mul_just_below_negative_1_rn_ftz_f() {
; CHECK-LABEL: define float @test_mul_just_below_negative_1_rn_ftz_f() {
; CHECK-NEXT: ret float -1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rn.ftz.f(float -0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float -0.75, float 0x3FF5555560000000, /* rnd=rn */ i32 1)
ret float %res
}
@@ -711,7 +1135,7 @@ define float @test_mul_just_below_negative_1_rp_ftz_f() {
; CHECK-LABEL: define float @test_mul_just_below_negative_1_rp_ftz_f() {
; CHECK-NEXT: ret float -1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rp.ftz.f(float -0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float -0.75, float 0x3FF5555560000000, /* rnd=rp */ i32 2)
ret float %res
}
@@ -719,7 +1143,7 @@ define float @test_mul_just_below_negative_1_rz_ftz_f() {
; CHECK-LABEL: define float @test_mul_just_below_negative_1_rz_ftz_f() {
; CHECK-NEXT: ret float -1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rz.ftz.f(float -0.75, float 0x3FF5555560000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float -0.75, float 0x3FF5555560000000, /* rnd=rz */ i32 0)
ret float %res
}
@@ -737,7 +1161,7 @@ define double @test_mul_just_below_negative_1_rm_d() {
; CHECK-LABEL: define double @test_mul_just_below_negative_1_rm_d() {
; CHECK-NEXT: ret double f0xBFF0000000000001
;
- %res = call double @llvm.nvvm.mul.rm.d(double -0.75, double 0x3FF5555555555556)
+ %res = call double @llvm.nvvm.fmul.f64(double -0.75, double 0x3FF5555555555556, /* rnd=rm */ i32 3)
ret double %res
}
@@ -745,7 +1169,7 @@ define double @test_mul_just_below_negative_1_rn_d() {
; CHECK-LABEL: define double @test_mul_just_below_negative_1_rn_d() {
; CHECK-NEXT: ret double -1.000000e+00
;
- %res = call double @llvm.nvvm.mul.rn.d(double -0.75, double 0x3FF5555555555556)
+ %res = call double @llvm.nvvm.fmul.f64(double -0.75, double 0x3FF5555555555556, /* rnd=rn */ i32 1)
ret double %res
}
@@ -753,7 +1177,7 @@ define double @test_mul_just_below_negative_1_rp_d() {
; CHECK-LABEL: define double @test_mul_just_below_negative_1_rp_d() {
; CHECK-NEXT: ret double -1.000000e+00
;
- %res = call double @llvm.nvvm.mul.rp.d(double -0.75, double 0x3FF5555555555556)
+ %res = call double @llvm.nvvm.fmul.f64(double -0.75, double 0x3FF5555555555556, /* rnd=rp */ i32 2)
ret double %res
}
@@ -761,10 +1185,178 @@ define double @test_mul_just_below_negative_1_rz_d() {
; CHECK-LABEL: define double @test_mul_just_below_negative_1_rz_d() {
; CHECK-NEXT: ret double -1.000000e+00
;
- %res = call double @llvm.nvvm.mul.rz.d(double -0.75, double 0x3FF5555555555556)
+ %res = call double @llvm.nvvm.fmul.f64(double -0.75, double 0x3FF5555555555556, /* rnd=rz */ i32 0)
ret double %res
}
+;###############################################################
+;# Mul(0.75, 4/3 + epsilon) #
+;###############################################################
+; Tests multiplication of 0.75 by a value slightly above 4/3,
+; where different rounding modes produce different results.
+; The exact result would be 1.0, but since 4/3 cannot be exactly encoded
+; as a half, the calculated result falls between 1.0 and 1.0 + 2^-10.
+; - RN, RZ, RM round to 1.0 (rounding to nearest/zero/down)
+; - RP rounds to 1.0 + 2^-10 (rounding up)
+
+define half @test_mul_just_above_1_rm_f16() {
+; CHECK-LABEL: define half @test_mul_just_above_1_rm_f16() {
+; CHECK-NEXT: ret half 1.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 0.75, half 0xH3D56, /* rnd=rm */ i32 3)
+ ret half %res
+}
+
+define half @test_mul_just_above_1_rn_f16() {
+; CHECK-LABEL: define half @test_mul_just_above_1_rn_f16() {
+; CHECK-NEXT: ret half 1.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 0.75, half 0xH3D56, /* rnd=rn */ i32 1)
+ ret half %res
+}
+
+define half @test_mul_just_above_1_rp_f16() {
+; CHECK-LABEL: define half @test_mul_just_above_1_rp_f16() {
+; CHECK-NEXT: ret half 1.000980e+00
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 0.75, half 0xH3D56, /* rnd=rp */ i32 2)
+ ret half %res
+}
+
+define half @test_mul_just_above_1_rz_f16() {
+; CHECK-LABEL: define half @test_mul_just_above_1_rz_f16() {
+; CHECK-NEXT: ret half 1.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.f16(half 0.75, half 0xH3D56, /* rnd=rz */ i32 0)
+ ret half %res
+}
+
+;###############################################################
+;# Mul(0.75, 4/3 + epsilon) #
+;###############################################################
+; Tests multiplication of 0.75 by a value slightly above 4/3,
+; where different rounding modes produce different results.
+; The exact result would be 1.0, but since 4/3 cannot be exactly encoded
+; as a bfloat, the calculated result falls between 1.0 and 1.0 + 2^-7.
+; - RN, RZ, RM round to 1.0 (rounding to nearest/zero/down)
+; - RP rounds to 1.0 + 2^-7 (rounding up)
+
+define bfloat @test_mul_just_above_1_rm_bf16() {
+; CHECK-LABEL: define bfloat @test_mul_just_above_1_rm_bf16() {
+; CHECK-NEXT: ret bfloat 1.000000e+00
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 0.75, bfloat 0xR3FAB, /* rnd=rm */ i32 3)
+ ret bfloat %res
+}
+
+define bfloat @test_mul_just_above_1_rn_bf16() {
+; CHECK-LABEL: define bfloat @test_mul_just_above_1_rn_bf16() {
+; CHECK-NEXT: ret bfloat 1.000000e+00
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 0.75, bfloat 0xR3FAB, /* rnd=rn */ i32 1)
+ ret bfloat %res
+}
+
+define bfloat @test_mul_just_above_1_rp_bf16() {
+; CHECK-LABEL: define bfloat @test_mul_just_above_1_rp_bf16() {
+; CHECK-NEXT: ret bfloat 1.007810e+00
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 0.75, bfloat 0xR3FAB, /* rnd=rp */ i32 2)
+ ret bfloat %res
+}
+
+define bfloat @test_mul_just_above_1_rz_bf16() {
+; CHECK-LABEL: define bfloat @test_mul_just_above_1_rz_bf16() {
+; CHECK-NEXT: ret bfloat 1.000000e+00
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat 0.75, bfloat 0xR3FAB, /* rnd=rz */ i32 0)
+ ret bfloat %res
+}
+
+;###############################################################
+;# Mul(-0.75, 4/3 + epsilon) #
+;###############################################################
+; Tests multiplication of -0.75 by a value slightly above 4/3,
+; where different rounding modes produce different results.
+; The exact result would be -1.0, but since 4/3 cannot be exactly encoded
+; as a half, the calculated result falls between -1.0 and -1.0 - 2^-10.
+; - RN, RZ, RP round to -1.0 (rounding to nearest/zero/up)
+; - RM rounds to -1.0 - 2^-10 (rounding down)
+
+define half @test_mul_just_below_negative_1_rm_f16() {
+; CHECK-LABEL: define half @test_mul_just_below_negative_1_rm_f16() {
+; CHECK-NEXT: ret half -1.000980e+00
+;
+ %res = call half @llvm.nvvm.fmul.f16(half -0.75, half 0xH3D56, /* rnd=rm */ i32 3)
+ ret half %res
+}
+
+define half @test_mul_just_below_negative_1_rn_f16() {
+; CHECK-LABEL: define half @test_mul_just_below_negative_1_rn_f16() {
+; CHECK-NEXT: ret half -1.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.f16(half -0.75, half 0xH3D56, /* rnd=rn */ i32 1)
+ ret half %res
+}
+
+define half @test_mul_just_below_negative_1_rp_f16() {
+; CHECK-LABEL: define half @test_mul_just_below_negative_1_rp_f16() {
+; CHECK-NEXT: ret half -1.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.f16(half -0.75, half 0xH3D56, /* rnd=rp */ i32 2)
+ ret half %res
+}
+
+define half @test_mul_just_below_negative_1_rz_f16() {
+; CHECK-LABEL: define half @test_mul_just_below_negative_1_rz_f16() {
+; CHECK-NEXT: ret half -1.000000e+00
+;
+ %res = call half @llvm.nvvm.fmul.f16(half -0.75, half 0xH3D56, /* rnd=rz */ i32 0)
+ ret half %res
+}
+
+;###############################################################
+;# Mul(-0.75, 4/3 + epsilon) #
+;###############################################################
+; Tests multiplication of -0.75 by a value slightly above 4/3,
+; where different rounding modes produce different results.
+; The exact result would be -1.0, but since 4/3 cannot be exactly encoded
+; as a bfloat, the calculated result falls between -1.0 and -1.0 - 2^-7.
+; - RN, RZ, RP round to -1.0 (rounding to nearest/zero/up)
+; - RM rounds to -1.0 - 2^-7 (rounding down)
+
+define bfloat @test_mul_just_below_negative_1_rm_bf16() {
+; CHECK-LABEL: define bfloat @test_mul_just_below_negative_1_rm_bf16() {
+; CHECK-NEXT: ret bfloat -1.007810e+00
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat -0.75, bfloat 0xR3FAB, /* rnd=rm */ i32 3)
+ ret bfloat %res
+}
+
+define bfloat @test_mul_just_below_negative_1_rn_bf16() {
+; CHECK-LABEL: define bfloat @test_mul_just_below_negative_1_rn_bf16() {
+; CHECK-NEXT: ret bfloat -1.000000e+00
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat -0.75, bfloat 0xR3FAB, /* rnd=rn */ i32 1)
+ ret bfloat %res
+}
+
+define bfloat @test_mul_just_below_negative_1_rp_bf16() {
+; CHECK-LABEL: define bfloat @test_mul_just_below_negative_1_rp_bf16() {
+; CHECK-NEXT: ret bfloat -1.000000e+00
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat -0.75, bfloat 0xR3FAB, /* rnd=rp */ i32 2)
+ ret bfloat %res
+}
+
+define bfloat @test_mul_just_below_negative_1_rz_bf16() {
+; CHECK-LABEL: define bfloat @test_mul_just_below_negative_1_rz_bf16() {
+; CHECK-NEXT: ret bfloat -1.000000e+00
+;
+ %res = call bfloat @llvm.nvvm.fmul.bf16(bfloat -0.75, bfloat 0xR3FAB, /* rnd=rz */ i32 0)
+ ret bfloat %res
+}
+
;###############################################################
;# Mul(0.625, 1.6 + epsilon) #
;###############################################################
@@ -778,7 +1370,7 @@ define float @test_mul_slightly_more_above_1_rm_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_above_1_rm_f() {
; CHECK-NEXT: ret float 1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rm.f(float 0x3FE4000000000000, float 0x3FF99999C0000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 0x3FE4000000000000, float 0x3FF99999C0000000, /* rnd=rm */ i32 3)
ret float %res
}
@@ -786,7 +1378,7 @@ define float @test_mul_slightly_more_above_1_rn_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_above_1_rn_f() {
; CHECK-NEXT: ret float f0x3F800001
;
- %res = call float @llvm.nvvm.mul.rn.f(float 0x3FE4000000000000, float 0x3FF99999C0000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 0x3FE4000000000000, float 0x3FF99999C0000000, /* rnd=rn */ i32 1)
ret float %res
}
@@ -794,7 +1386,7 @@ define float @test_mul_slightly_more_above_1_rp_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_above_1_rp_f() {
; CHECK-NEXT: ret float f0x3F800001
;
- %res = call float @llvm.nvvm.mul.rp.f(float 0x3FE4000000000000, float 0x3FF99999C0000000 )
+ %res = call float @llvm.nvvm.fmul.f32(float 0x3FE4000000000000, float 0x3FF99999C0000000, /* rnd=rp */ i32 2)
ret float %res
}
@@ -802,7 +1394,7 @@ define float @test_mul_slightly_more_above_1_rz_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_above_1_rz_f() {
; CHECK-NEXT: ret float 1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rz.f(float 0x3FE4000000000000, float 0x3FF99999C0000000 )
+ %res = call float @llvm.nvvm.fmul.f32(float 0x3FE4000000000000, float 0x3FF99999C0000000, /* rnd=rz */ i32 0)
ret float %res
}
@@ -810,7 +1402,7 @@ define float @test_mul_slightly_more_above_1_rm_ftz_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_above_1_rm_ftz_f() {
; CHECK-NEXT: ret float 1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rm.ftz.f(float 0x3FE4000000000000, float 0x3FF99999C0000000 )
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0x3FE4000000000000, float 0x3FF99999C0000000, /* rnd=rm */ i32 3)
ret float %res
}
@@ -818,7 +1410,7 @@ define float @test_mul_slightly_more_above_1_rn_ftz_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_above_1_rn_ftz_f() {
; CHECK-NEXT: ret float f0x3F800001
;
- %res = call float @llvm.nvvm.mul.rn.ftz.f(float 0x3FE4000000000000, float 0x3FF99999C0000000 )
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0x3FE4000000000000, float 0x3FF99999C0000000, /* rnd=rn */ i32 1)
ret float %res
}
@@ -826,7 +1418,7 @@ define float @test_mul_slightly_more_above_1_rp_ftz_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_above_1_rp_ftz_f() {
; CHECK-NEXT: ret float f0x3F800001
;
- %res = call float @llvm.nvvm.mul.rp.ftz.f(float 0x3FE4000000000000, float 0x3FF99999C0000000 )
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0x3FE4000000000000, float 0x3FF99999C0000000, /* rnd=rp */ i32 2)
ret float %res
}
@@ -834,7 +1426,7 @@ define float @test_mul_slightly_more_above_1_rz_ftz_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_above_1_rz_ftz_f() {
; CHECK-NEXT: ret float 1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rz.ftz.f(float 0x3FE4000000000000, float 0x3FF99999C0000000 )
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0x3FE4000000000000, float 0x3FF99999C0000000, /* rnd=rz */ i32 0)
ret float %res
}
@@ -851,7 +1443,7 @@ define double @test_mul_slightly_more_above_1_rm_d() {
; CHECK-LABEL: define double @test_mul_slightly_more_above_1_rm_d() {
; CHECK-NEXT: ret double 1.000000e+00
;
- %res = call double @llvm.nvvm.mul.rm.d(double 0x3FE4000000000000, double 0x3FF999999999999B)
+ %res = call double @llvm.nvvm.fmul.f64(double 0x3FE4000000000000, double 0x3FF999999999999B, /* rnd=rm */ i32 3)
ret double %res
}
@@ -859,7 +1451,7 @@ define double @test_mul_slightly_more_above_1_rn_d() {
; CHECK-LABEL: define double @test_mul_slightly_more_above_1_rn_d() {
; CHECK-NEXT: ret double f0x3FF0000000000001
;
- %res = call double @llvm.nvvm.mul.rn.d(double 0x3FE4000000000000, double 0x3FF999999999999B)
+ %res = call double @llvm.nvvm.fmul.f64(double 0x3FE4000000000000, double 0x3FF999999999999B, /* rnd=rn */ i32 1)
ret double %res
}
@@ -867,7 +1459,7 @@ define double @test_mul_slightly_more_above_1_rp_d() {
; CHECK-LABEL: define double @test_mul_slightly_more_above_1_rp_d() {
; CHECK-NEXT: ret double f0x3FF0000000000001
;
- %res = call double @llvm.nvvm.mul.rp.d(double 0x3FE4000000000000, double 0x3FF999999999999B)
+ %res = call double @llvm.nvvm.fmul.f64(double 0x3FE4000000000000, double 0x3FF999999999999B, /* rnd=rp */ i32 2)
ret double %res
}
@@ -875,7 +1467,7 @@ define double @test_mul_slightly_more_above_1_rz_d() {
; CHECK-LABEL: define double @test_mul_slightly_more_above_1_rz_d() {
; CHECK-NEXT: ret double 1.000000e+00
;
- %res = call double @llvm.nvvm.mul.rz.d(double 0x3FE4000000000000, double 0x3FF999999999999B)
+ %res = call double @llvm.nvvm.fmul.f64(double 0x3FE4000000000000, double 0x3FF999999999999B, /* rnd=rz */ i32 0)
ret double %res
}
@@ -892,7 +1484,7 @@ define float @test_mul_slightly_more_below_negative_1_rm_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_below_negative_1_rm_f() {
; CHECK-NEXT: ret float f0xBF800001
;
- %res = call float @llvm.nvvm.mul.rm.f(float 0x3FE4000000000000, float 0xBFF99999C0000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 0x3FE4000000000000, float 0xBFF99999C0000000, /* rnd=rm */ i32 3)
ret float %res
}
@@ -900,7 +1492,7 @@ define float @test_mul_slightly_more_below_negative_1_rn_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_below_negative_1_rn_f() {
; CHECK-NEXT: ret float f0xBF800001
;
- %res = call float @llvm.nvvm.mul.rn.f(float 0x3FE4000000000000, float 0xBFF99999C0000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 0x3FE4000000000000, float 0xBFF99999C0000000, /* rnd=rn */ i32 1)
ret float %res
}
@@ -908,7 +1500,7 @@ define float @test_mul_slightly_more_below_negative_1_rp_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_below_negative_1_rp_f() {
; CHECK-NEXT: ret float -1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rp.f(float 0x3FE4000000000000, float 0xBFF99999C0000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 0x3FE4000000000000, float 0xBFF99999C0000000, /* rnd=rp */ i32 2)
ret float %res
}
@@ -916,7 +1508,7 @@ define float @test_mul_slightly_more_below_negative_1_rz_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_below_negative_1_rz_f() {
; CHECK-NEXT: ret float -1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rz.f(float 0x3FE4000000000000, float 0xBFF99999C0000000)
+ %res = call float @llvm.nvvm.fmul.f32(float 0x3FE4000000000000, float 0xBFF99999C0000000, /* rnd=rz */ i32 0)
ret float %res
}
@@ -924,7 +1516,7 @@ define float @test_mul_slightly_more_below_negative_1_rm_ftz_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_below_negative_1_rm_ftz_f() {
; CHECK-NEXT: ret float f0xBF800001
;
- %res = call float @llvm.nvvm.mul.rm.ftz.f(float 0x3FE4000000000000, float 0xBFF99999C0000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0x3FE4000000000000, float 0xBFF99999C0000000, /* rnd=rm */ i32 3)
ret float %res
}
@@ -932,7 +1524,7 @@ define float @test_mul_slightly_more_below_negative_1_rn_ftz_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_below_negative_1_rn_ftz_f() {
; CHECK-NEXT: ret float f0xBF800001
;
- %res = call float @llvm.nvvm.mul.rn.ftz.f(float 0x3FE4000000000000, float 0xBFF99999C0000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0x3FE4000000000000, float 0xBFF99999C0000000, /* rnd=rn */ i32 1)
ret float %res
}
@@ -940,7 +1532,7 @@ define float @test_mul_slightly_more_below_negative_1_rp_ftz_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_below_negative_1_rp_ftz_f() {
; CHECK-NEXT: ret float -1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rp.ftz.f(float 0x3FE4000000000000, float 0xBFF99999C0000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0x3FE4000000000000, float 0xBFF99999C0000000, /* rnd=rp */ i32 2)
ret float %res
}
@@ -948,7 +1540,7 @@ define float @test_mul_slightly_more_below_negative_1_rz_ftz_f() {
; CHECK-LABEL: define float @test_mul_slightly_more_below_negative_1_rz_ftz_f() {
; CHECK-NEXT: ret float -1.000000e+00
;
- %res = call float @llvm.nvvm.mul.rz.ftz.f(float 0x3FE4000000000000, float 0xBFF99999C0000000)
+ %res = call float @llvm.nvvm.fmul.ftz.f32(float 0x3FE4000000000000, float 0xBFF99999C0000000, /* rnd=rz */ i32 0)
ret float %res
}
@@ -965,7 +1557,7 @@ define double @test_mul_slightly_more_below_negative_1_rm_d() {
; CHECK-LABEL: define double @test_mul_slightly_more_below_negative_1_rm_d() {
; CHECK-NEXT: ret double f0xBFF0000000000001
;
- %res = call double @llvm.nvvm.mul.rm.d(double 0x3FE4000000000000, double 0xBFF999999999999B)
+ %res = call double @llvm.nvvm.fmul.f64(double 0x3FE4000000000000, double 0xBFF999999999999B, /* rnd=rm */ i32 3)
ret double %res
}
@@ -973,7 +1565,7 @@ define double @test_mul_slightly_more_below_negative_1_rn_d() {
; CHECK-LABEL: define double @test_mul_slightly_more_below_negative_1_rn_d() {
; CHECK-NEXT: ret double f0xBFF0000000000001
;
- %res = call double @llvm.nvvm.mul.rn.d(double 0x3FE4000000000000, double 0xBFF999999999999B)
+ %res = call double @llvm.nvvm.fmul.f64(double 0x3FE4000000000000, double 0xBFF999999999999B, /* rnd=rn */ i32 1)
ret double %res
}
@@ -981,7 +1573,7 @@ define double @test_mul_slightly_more_below_negative_1_rp_d() {
; CHECK-LABEL: define double @test_mul_slightly_more_below_negative_1_rp_d() {
; CHECK-NEXT: ret double -1.000000e+00
;
- %res = call double @llvm.nvvm.mul.rp.d(double 0x3FE4000000000000, double 0xBFF999999999999B)
+ %res = call double @llvm.nvvm.fmul.f64(double 0x3FE4000000000000, double 0xBFF999999999999B, /* rnd=rp */ i32 2)
ret double %res
}
@@ -989,6 +1581,6 @@ define double @test_mul_slightly_more_below_negative_1_rz_d() {
; CHECK-LABEL: define double @test_mul_slightly_more_below_negative_1_rz_d() {
; CHECK-NEXT: ret double -1.000000e+00
;
- %res = call double @llvm.nvvm.mul.rz.d(double 0x3FE4000000000000, double 0xBFF999999999999B)
+ %res = call double @llvm.nvvm.fmul.f64(double 0x3FE4000000000000, double 0xBFF999999999999B, /* rnd=rz */ i32 0)
ret double %res
}
diff --git a/llvm/test/Verifier/NVPTX/fmul.ll b/llvm/test/Verifier/NVPTX/fmul.ll
new file mode 100644
index 00000000000000..2e5779466b4b11
--- /dev/null
+++ b/llvm/test/Verifier/NVPTX/fmul.ll
@@ -0,0 +1,16 @@
+; RUN: not llvm-as %s -o /dev/null 2>&1 | FileCheck %s
+
+declare float @llvm.nvvm.fmul.f32(float, float, i32 immarg)
+
+define void @test_fmul_rounding_mode(float %a) {
+ ; CHECK: immarg value 4 for arg 2 out of range [0,4)
+ call float @llvm.nvvm.fmul.f32(float %a, float %a, i32 4)
+
+ ; CHECK: immarg value 7 for arg 2 out of range [0,4)
+ call float @llvm.nvvm.fmul.f32(float %a, float %a, i32 7)
+
+ ; CHECK: immarg value -1 for arg 2 out of range [0,4)
+ call float @llvm.nvvm.fmul.f32(float %a, float %a, i32 -1)
+
+ ret void
+}
diff --git a/llvm/test/Verifier/intrinsic-bad-arg-type1.ll b/llvm/test/Verifier/intrinsic-bad-arg-type1.ll
index 05d9605886d394..f9b02ada52cbc5 100644
--- a/llvm/test/Verifier/intrinsic-bad-arg-type1.ll
+++ b/llvm/test/Verifier/intrinsic-bad-arg-type1.ll
@@ -28,8 +28,8 @@ declare i32 @llvm.call.preallocated.setup(i32)
declare double @llvm.fptrunc.round.f64.f64(double, i16)
; CHECK: intrinsic argument 0 type expected half, but got i16
-; CHECK-NEXT: declare half @llvm.nvvm.mul.rn.sat.f16(i16, half)
-declare half @llvm.nvvm.mul.rn.sat.f16(i16, half)
+; CHECK-NEXT: declare half @llvm.nvvm.fma.rn.f16(i16, half, half)
+declare half @llvm.nvvm.fma.rn.f16(i16, half, half)
; CHECK: intrinsic return type expected bfloat, but got half
; CHECK-NEXT: declare half @llvm.arm.neon.vcvtbfp2bf(float)
diff --git a/llvm/unittests/IR/IntrinsicsTest.cpp b/llvm/unittests/IR/IntrinsicsTest.cpp
index f982715154bc05..ea4d44d6e2f274 100644
--- a/llvm/unittests/IR/IntrinsicsTest.cpp
+++ b/llvm/unittests/IR/IntrinsicsTest.cpp
@@ -118,7 +118,7 @@ TEST(IntrinsicNameLookup, ClangBuiltinLookup) {
{"__builtin_HEXAGON_A2_tfr", "hexagon", hexagon_A2_tfr},
{"__builtin_lasx_xbz_w", "loongarch", loongarch_lasx_xbz_w},
{"__builtin_mips_bitrev", "mips", mips_bitrev},
- {"__nvvm_mul_rn_d", "nvvm", nvvm_mul_rn_d},
+ {"__nvvm_div_rn_d", "nvvm", nvvm_div_rn_d},
{"__builtin_altivec_dss", "ppc", ppc_altivec_dss},
{"__builtin_riscv_sha512sum1r", "riscv", riscv_sha512sum1r},
{"__builtin_tend", "s390", s390_tend},
>From 13bd7ec0cd6064b512b9f8baddaf528b8011124e Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Tue, 22 Sep 2026 07:27:17 +0000
Subject: [PATCH 2/2] address comments
---
llvm/docs/NVPTXUsage.md | 60 ++++++++++++---------
llvm/include/llvm/IR/IntrinsicsNVVM.td | 7 +--
llvm/include/llvm/IR/NVVMIntrinsicUtils.h | 2 +-
llvm/lib/Analysis/ConstantFolding.cpp | 2 +-
llvm/lib/IR/AutoUpgrade.cpp | 6 +--
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 13 +++--
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 12 ++---
7 files changed, 53 insertions(+), 49 deletions(-)
diff --git a/llvm/docs/NVPTXUsage.md b/llvm/docs/NVPTXUsage.md
index 1386046f7b85d2..8ef068a1f23ece 100644
--- a/llvm/docs/NVPTXUsage.md
+++ b/llvm/docs/NVPTXUsage.md
@@ -1347,28 +1347,34 @@ vectors is added to `%c` to produce the return.
##### Syntax:
-This is an overloaded intrinsic. The '`.ftz`' and '`.sat`' modifiers are
-optional.
+This is an overloaded intrinsic of the form:
```llvm
-declare half @llvm.nvvm.fadd{.ftz}{.sat}.f16(half %a, half %b, i32 immarg %rnd)
-declare <2 x half> @llvm.nvvm.fadd{.ftz}{.sat}.v2f16(<2 x half> %a, <2 x half> %b, i32 immarg %rnd)
-declare bfloat @llvm.nvvm.fadd.bf16(bfloat %a, bfloat %b, i32 immarg %rnd)
-declare <2 x bfloat> @llvm.nvvm.fadd.v2bf16(<2 x bfloat> %a, <2 x bfloat> %b, i32 immarg %rnd)
-declare float @llvm.nvvm.fadd{.ftz}{.sat}.f32(float %a, float %b, i32 immarg %rnd)
-declare <2 x float> @llvm.nvvm.fadd{.ftz}.v2f32(<2 x float> %a, <2 x float> %b, i32 immarg %rnd)
-declare double @llvm.nvvm.fadd.f64(double %a, double %b, i32 immarg %rnd)
+declare <ty> @llvm.nvvm.fadd{.ftz}{.sat}.<sfx>(<ty> %a, <ty> %b, i32 immarg %flag_fp_rnd_mode)
+```
+
+where '`<sfx>`' is the mangled suffix of the overloaded type '`<ty>`' and the
+'`.ftz`' and '`.sat`' modifiers are optional. The supported variants are:
+
+```llvm
+declare half @llvm.nvvm.fadd{.ftz}{.sat}.f16(half %a, half %b, i32 immarg %flag_fp_rnd_mode)
+declare <2 x half> @llvm.nvvm.fadd{.ftz}{.sat}.v2f16(<2 x half> %a, <2 x half> %b, i32 immarg %flag_fp_rnd_mode)
+declare bfloat @llvm.nvvm.fadd.bf16(bfloat %a, bfloat %b, i32 immarg %flag_fp_rnd_mode)
+declare <2 x bfloat> @llvm.nvvm.fadd.v2bf16(<2 x bfloat> %a, <2 x bfloat> %b, i32 immarg %flag_fp_rnd_mode)
+declare float @llvm.nvvm.fadd{.ftz}{.sat}.f32(float %a, float %b, i32 immarg %flag_fp_rnd_mode)
+declare <2 x float> @llvm.nvvm.fadd{.ftz}.v2f32(<2 x float> %a, <2 x float> %b, i32 immarg %flag_fp_rnd_mode)
+declare double @llvm.nvvm.fadd.f64(double %a, double %b, i32 immarg %flag_fp_rnd_mode)
```
##### Overview:
The '`llvm.nvvm.fadd.*`' intrinsics add `%a` and `%b` using the rounding mode
-selected by `%rnd` and the modifiers present in the intrinsic name. They
-correspond directly to the `add` PTX instruction.
+selected by `%flag_fp_rnd_mode` and the modifiers present in the intrinsic
+name. They correspond directly to the `add` PTX instruction.
##### Semantics:
-`%rnd` selects the rounding mode applied to the result, see
+`%flag_fp_rnd_mode` selects the rounding mode applied to the result, see
{ref}`fp-rounding-modes`.
The '`.ftz`' modifier flushes subnormal inputs and results to sign-preserving
@@ -1407,28 +1413,34 @@ PTX instruction. The supported combinations are:
##### Syntax:
-This is an overloaded intrinsic. The '`.ftz`' and '`.sat`' modifiers are
-optional.
+This is an overloaded intrinsic of the form:
+
+```llvm
+declare <ty> @llvm.nvvm.fmul{.ftz}{.sat}.<sfx>(<ty> %a, <ty> %b, i32 immarg %flag_fp_rnd_mode)
+```
+
+where '`<sfx>`' is the mangled suffix of the overloaded type '`<ty>`' and the
+'`.ftz`' and '`.sat`' modifiers are optional. The supported variants are:
```llvm
-declare half @llvm.nvvm.fmul{.ftz}{.sat}.f16(half %a, half %b, i32 immarg %rnd)
-declare <2 x half> @llvm.nvvm.fmul{.ftz}{.sat}.v2f16(<2 x half> %a, <2 x half> %b, i32 immarg %rnd)
-declare bfloat @llvm.nvvm.fmul.bf16(bfloat %a, bfloat %b, i32 immarg %rnd)
-declare <2 x bfloat> @llvm.nvvm.fmul.v2bf16(<2 x bfloat> %a, <2 x bfloat> %b, i32 immarg %rnd)
-declare float @llvm.nvvm.fmul{.ftz}{.sat}.f32(float %a, float %b, i32 immarg %rnd)
-declare <2 x float> @llvm.nvvm.fmul{.ftz}.v2f32(<2 x float> %a, <2 x float> %b, i32 immarg %rnd)
-declare double @llvm.nvvm.fmul.f64(double %a, double %b, i32 immarg %rnd)
+declare half @llvm.nvvm.fmul{.ftz}{.sat}.f16(half %a, half %b, i32 immarg %flag_fp_rnd_mode)
+declare <2 x half> @llvm.nvvm.fmul{.ftz}{.sat}.v2f16(<2 x half> %a, <2 x half> %b, i32 immarg %flag_fp_rnd_mode)
+declare bfloat @llvm.nvvm.fmul.bf16(bfloat %a, bfloat %b, i32 immarg %flag_fp_rnd_mode)
+declare <2 x bfloat> @llvm.nvvm.fmul.v2bf16(<2 x bfloat> %a, <2 x bfloat> %b, i32 immarg %flag_fp_rnd_mode)
+declare float @llvm.nvvm.fmul{.ftz}{.sat}.f32(float %a, float %b, i32 immarg %flag_fp_rnd_mode)
+declare <2 x float> @llvm.nvvm.fmul{.ftz}.v2f32(<2 x float> %a, <2 x float> %b, i32 immarg %flag_fp_rnd_mode)
+declare double @llvm.nvvm.fmul.f64(double %a, double %b, i32 immarg %flag_fp_rnd_mode)
```
##### Overview:
The '`llvm.nvvm.fmul.*`' intrinsics multiply `%a` and `%b` using the rounding
-mode selected by `%rnd` and the modifiers present in the intrinsic name. They
-correspond directly to the `mul` PTX instruction.
+mode selected by `%flag_fp_rnd_mode` and the modifiers present in the
+intrinsic name. They correspond directly to the `mul` PTX instruction.
##### Semantics:
-`%rnd` selects the rounding mode applied to the result, see
+`%flag_fp_rnd_mode` selects the rounding mode applied to the result, see
{ref}`fp-rounding-modes`.
The '`.ftz`' modifier flushes subnormal inputs and results to sign-preserving
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 50d32c305571b6..259151b687e0b3 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1683,15 +1683,12 @@ let TargetPrefix = "nvvm" in {
let IntrProperties = [IntrNoMem, IntrSpeculatable, Commutative,
IntrNoCreateUndefOrPoison, ImmArg<ArgIndex<2>>,
Range<ArgIndex<2>, 0, 4>,
- ArgInfo<ArgIndex<2>,
- [ArgName<"rnd">,
- ImmArgPrinter<"printFPRoundingMode">]>] in
+ ArgInfo<ArgIndex<2>, [ArgName<"rnd">, ImmArgPrinter<"printFPRoundingMode">]>] in
foreach ftz = ["", "_ftz"] in
foreach sat = ["", "_sat"] in
def int_nvvm_fmul # ftz # sat :
DefaultAttrsIntrinsic<[llvm_anyfloat_ty],
- [LLVMMatchType<0>, LLVMMatchType<0>,
- llvm_i32_ty]>;
+ [LLVMMatchType<0>, LLVMMatchType<0>, llvm_i32_ty]>;
//
// Div
diff --git a/llvm/include/llvm/IR/NVVMIntrinsicUtils.h b/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
index 40efa898d27281..c92fa1cf21f1bb 100644
--- a/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
+++ b/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
@@ -714,7 +714,7 @@ inline bool FPArithShouldFTZ(Intrinsic::ID IntrinsicID) {
llvm_unreachable("Checking FTZ flag for invalid NVVM fadd/fmul intrinsic");
}
-inline bool FPArithIsSat(Intrinsic::ID IntrinsicID) {
+inline bool FPArithIsSaturating(Intrinsic::ID IntrinsicID) {
switch (IntrinsicID) {
case Intrinsic::nvvm_fadd_sat:
case Intrinsic::nvvm_fadd_ftz_sat:
diff --git a/llvm/lib/Analysis/ConstantFolding.cpp b/llvm/lib/Analysis/ConstantFolding.cpp
index bca11a4dac0da6..b0c07c986c8a51 100644
--- a/llvm/lib/Analysis/ConstantFolding.cpp
+++ b/llvm/lib/Analysis/ConstantFolding.cpp
@@ -1997,7 +1997,7 @@ static bool canConstantFoldIntrinsic(Intrinsic::ID ID, bool IsStrictFP) {
case Intrinsic::nvvm_sqrt_rn_ftz_f:
return !IsStrictFP;
- // NVVM add/mul intrinsics with explicit rounding modes
+ // NVVM fadd/fmul intrinsics with explicit rounding modes
case Intrinsic::nvvm_fadd:
case Intrinsic::nvvm_fadd_ftz:
case Intrinsic::nvvm_fmul:
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index ad0ccdf6396c36..b11d8445b93f3e 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -3238,9 +3238,9 @@ void llvm::UpgradeInlineAsmString(std::string *AsmStr) {
static Value *upgradeNVVMFPArithCall(IRBuilder<> &Builder, CallBase *CI,
StringRef Name,
const Intrinsic::ID IIDs[2][2]) {
- auto Upgrade = getNVVMFPArithUpgrade(Name, IIDs);
- assert(Upgrade && "unsupported nvvm.add.*/nvvm.mul.* intrinsic");
- auto [IID, RoundingMode] = *Upgrade;
+ auto Result = getNVVMFPArithUpgrade(Name, IIDs);
+ assert(Result && "unsupported nvvm.add.*/nvvm.mul.* intrinsic");
+ auto [IID, RoundingMode] = *Result;
Value *A = CI->getArgOperand(0);
return Builder.CreateIntrinsic(
A->getType(), IID,
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 3ad43d336f1760..d624dbe13d31f0 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -7179,7 +7179,7 @@ static SDValue sinkProxyReg(SDValue R, SDValue Chain,
static unsigned getFAddWithNegOpcode(EVT VT, Intrinsic::ID IID,
APFloat::roundingMode RoundingMode) {
const bool IsFTZ = nvvm::FPArithShouldFTZ(IID);
- const bool IsSat = nvvm::FPArithIsSat(IID);
+ const bool IsSat = nvvm::FPArithIsSaturating(IID);
switch (VT.getScalarType().getSimpleVT().SimpleTy) {
case MVT::f16: {
static constexpr unsigned SubRNOpcodes[2][2] = {
@@ -7234,15 +7234,16 @@ static SDValue combineFAddWithNeg(SDNode *N, SelectionDAG &DAG,
// TODO: Remove the type-legality checks here once
// https://github.com/llvm/llvm-project/pull/172442 lands, adding support for
// explicit type constraints for overloaded intrinsics in tablegen.
-static bool isSupportedFPArith(EVT VT, const NVPTXSubtarget &STI,
+static bool isSupportedFPArith(SDNode *N, const NVPTXSubtarget &STI,
unsigned ISDOpcode, Intrinsic::ID IID,
APFloat::roundingMode RoundingMode) {
+ const EVT VT = N->getValueType(0);
if (VT.isVector() && VT.getVectorElementCount() != ElementCount::getFixed(2))
return false;
const bool IsRN = RoundingMode == APFloat::rmNearestTiesToEven;
const bool IsFTZ = nvvm::FPArithShouldFTZ(IID);
- const bool IsSat = nvvm::FPArithIsSat(IID);
+ const bool IsSat = nvvm::FPArithIsSaturating(IID);
switch (VT.getScalarType().getSimpleVT().SimpleTy) {
case MVT::f16:
return IsRN;
@@ -7285,8 +7286,7 @@ static SDValue combineIntrinsicWOChain(SDNode *N,
case Intrinsic::nvvm_fadd_ftz_sat: {
const auto RoundingMode = static_cast<APFloat::roundingMode>(
N->getConstantOperandAPInt(3).getSExtValue());
- if (!isSupportedFPArith(N->getValueType(0), STI, ISD::FADD, IID,
- RoundingMode))
+ if (!isSupportedFPArith(N, STI, ISD::FADD, IID, RoundingMode))
return diagnoseUnsupportedFPArith(N, DCI.DAG, IID, RoundingMode);
return combineFAddWithNeg(N, DCI.DAG, IID, RoundingMode);
}
@@ -7296,8 +7296,7 @@ static SDValue combineIntrinsicWOChain(SDNode *N,
case Intrinsic::nvvm_fmul_ftz_sat: {
const auto RoundingMode = static_cast<APFloat::roundingMode>(
N->getConstantOperandAPInt(3).getSExtValue());
- if (!isSupportedFPArith(N->getValueType(0), STI, ISD::FMUL, IID,
- RoundingMode))
+ if (!isSupportedFPArith(N, STI, ISD::FMUL, IID, RoundingMode))
return diagnoseUnsupportedFPArith(N, DCI.DAG, IID, RoundingMode);
break;
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index c245997e7c2f45..c7357b5b7de3dd 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2192,8 +2192,7 @@ foreach t = [F16RT, F16X2RT] in
foreach sat = ["", "sat"] in
def INT_NVVM_MUL_RN_ # !toupper(StrJoin<"_", [ftz, sat, t.PtxType]>.ret) :
F_MATH_2_RND_TY<StrJoin<".", ["mul.rn", ftz, sat, t.PtxType]>.ret, t,
- !cast<Intrinsic>(
- StrJoin<"_", ["int_nvvm_fmul", ftz, sat]>.ret),
+ !cast<Intrinsic>(StrJoin<"_", ["int_nvvm_fmul", ftz, sat]>.ret),
fp_rnd_rn_imm>;
// bf16/bf16x2
@@ -2206,15 +2205,12 @@ foreach t = [BF16RT, BF16X2RT] in
foreach ftz = ["", "ftz"] in {
foreach sat = ["", "sat"] in
def StrJoin<"_", ["INT_NVVM_MUL", !toupper(ftz), !toupper(sat), "F"]>.ret :
- F_MATH_2_RNDOP_TY<StrJoin<".", ["mul.${rnd}", ftz, sat, "f32"]>.ret,
- F32RT,
- !cast<Intrinsic>(
- StrJoin<"_", ["int_nvvm_fmul", ftz, sat]>.ret)>;
+ F_MATH_2_RNDOP_TY<StrJoin<".", ["mul.${rnd}", ftz, sat, "f32"]>.ret, F32RT,
+ !cast<Intrinsic>(StrJoin<"_", ["int_nvvm_fmul", ftz, sat]>.ret)>;
def StrJoin<"_", ["INT_NVVM_MUL", !toupper(ftz), "F32X2"]>.ret :
F_MATH_2_RNDOP_TY<StrJoin<".", ["mul.${rnd}", ftz, "f32x2"]>.ret, F32X2RT,
- !cast<Intrinsic>(
- StrJoin<"_", ["int_nvvm_fmul", ftz]>.ret),
+ !cast<Intrinsic>(StrJoin<"_", ["int_nvvm_fmul", ftz]>.ret),
[hasF32x2Instructions]>;
}
More information about the llvm-commits
mailing list