[clang] [llvm] Clang: emit llvm.minnum and llvm.maxnum with nsz always (PR #113133)
YunQiang Su via cfe-commits
cfe-commits at lists.llvm.org
Mon Oct 21 18:56:23 PDT 2024
https://github.com/wzssyqa updated https://github.com/llvm/llvm-project/pull/113133
>From ec7d885a7bc89fcb35c08890fde2f07da1e61984 Mon Sep 17 00:00:00 2001
From: YunQiang Su <yunqiang at isrc.iscas.ac.cn>
Date: Mon, 21 Oct 2024 15:18:38 +0800
Subject: [PATCH 1/2] Clang: emit llvm.minnum and llvm.maxnum with nsz always
See: https://github.com/llvm/llvm-project/pull/112852
We will define llvm.minnum and llvm.maxnum with +0.0>-0.0, by default,
while libc doesn't require it.
---
clang/lib/CodeGen/CGBuiltin.cpp | 33 ++++++++++++--------
clang/test/CodeGen/fmaxnum_fminnum_use_nsz.c | 33 ++++++++++++++++++++
llvm/include/llvm/IR/IRBuilder.h | 16 ++++++----
llvm/lib/IR/IRBuilder.cpp | 4 +--
4 files changed, 65 insertions(+), 21 deletions(-)
create mode 100644 clang/test/CodeGen/fmaxnum_fminnum_use_nsz.c
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 28f28c70b5ae52..f2d6049908720b 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -510,19 +510,20 @@ static Value *emitUnaryMaybeConstrainedFPBuiltin(CodeGenFunction &CGF,
// Emit an intrinsic that has 2 operands of the same type as its result.
// Depending on mode, this may be a constrained floating-point intrinsic.
-static Value *emitBinaryMaybeConstrainedFPBuiltin(CodeGenFunction &CGF,
- const CallExpr *E, unsigned IntrinsicID,
- unsigned ConstrainedIntrinsicID) {
+static Value *emitBinaryMaybeConstrainedFPBuiltin(
+ CodeGenFunction &CGF, const CallExpr *E, unsigned IntrinsicID,
+ unsigned ConstrainedIntrinsicID, llvm::FastMathFlags *FMF = nullptr) {
llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));
llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1));
CodeGenFunction::CGFPOptionsRAII FPOptsRAII(CGF, E);
if (CGF.Builder.getIsFPConstrained()) {
Function *F = CGF.CGM.getIntrinsic(ConstrainedIntrinsicID, Src0->getType());
- return CGF.Builder.CreateConstrainedFPCall(F, { Src0, Src1 });
+ return CGF.Builder.CreateConstrainedFPCall(F, {Src0, Src1}, "",
+ std::nullopt, std::nullopt, FMF);
} else {
Function *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType());
- return CGF.Builder.CreateCall(F, { Src0, Src1 });
+ return CGF.Builder.CreateCall(F, {Src0, Src1}, "", nullptr, FMF);
}
}
@@ -2846,10 +2847,13 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
case Builtin::BI__builtin_fmaxf:
case Builtin::BI__builtin_fmaxf16:
case Builtin::BI__builtin_fmaxl:
- case Builtin::BI__builtin_fmaxf128:
- return RValue::get(emitBinaryMaybeConstrainedFPBuiltin(*this, E,
- Intrinsic::maxnum,
- Intrinsic::experimental_constrained_maxnum));
+ case Builtin::BI__builtin_fmaxf128: {
+ llvm::FastMathFlags FMF;
+ FMF.setNoSignedZeros();
+ return RValue::get(emitBinaryMaybeConstrainedFPBuiltin(
+ *this, E, Intrinsic::maxnum,
+ Intrinsic::experimental_constrained_maxnum, &FMF));
+ }
case Builtin::BIfmin:
case Builtin::BIfminf:
@@ -2858,10 +2862,13 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
case Builtin::BI__builtin_fminf:
case Builtin::BI__builtin_fminf16:
case Builtin::BI__builtin_fminl:
- case Builtin::BI__builtin_fminf128:
- return RValue::get(emitBinaryMaybeConstrainedFPBuiltin(*this, E,
- Intrinsic::minnum,
- Intrinsic::experimental_constrained_minnum));
+ case Builtin::BI__builtin_fminf128: {
+ llvm::FastMathFlags FMF;
+ FMF.setNoSignedZeros();
+ return RValue::get(emitBinaryMaybeConstrainedFPBuiltin(
+ *this, E, Intrinsic::minnum,
+ Intrinsic::experimental_constrained_minnum, &FMF));
+ }
case Builtin::BIfmaximum_num:
case Builtin::BIfmaximum_numf:
diff --git a/clang/test/CodeGen/fmaxnum_fminnum_use_nsz.c b/clang/test/CodeGen/fmaxnum_fminnum_use_nsz.c
new file mode 100644
index 00000000000000..9798baf0432fea
--- /dev/null
+++ b/clang/test/CodeGen/fmaxnum_fminnum_use_nsz.c
@@ -0,0 +1,33 @@
+// RUN: %clang_cc1 -triple x86_64 %s -emit-llvm -o - 2>&1 | FileCheck %s --check-prefix=CHECK
+
+float fminf (float, float);
+double fmin (double, double);
+long double fminl (long double, long double);
+float fmaxf (float, float);
+double fmax (double, double);
+long double fmaxl (long double, long double);
+
+// CHECK: call nsz float @llvm.minnum.f32
+float fmin1(float a, float b) {
+ return fminf(a, b);
+}
+// CHECK: call nsz double @llvm.minnum.f64
+float fmin2(double a, double b) {
+ return fmin(a, b);
+}
+// CHECK: call nsz x86_fp80 @llvm.minnum.f80
+float fmin3(long double a, long double b) {
+ return fminl(a, b);
+}
+// CHECK: call nsz float @llvm.maxnum.f32
+float fmax1(float a, float b) {
+ return fmaxf(a, b);
+}
+// CHECK: call nsz double @llvm.maxnum.f64
+float fmax2(double a, double b) {
+ return fmax(a, b);
+}
+// CHECK: call nsz x86_fp80 @llvm.maxnum.f80
+float fmax3(long double a, long double b) {
+ return fmaxl(a, b);
+}
diff --git a/llvm/include/llvm/IR/IRBuilder.h b/llvm/include/llvm/IR/IRBuilder.h
index 23fd8350a29b3d..1baca4f003cad6 100644
--- a/llvm/include/llvm/IR/IRBuilder.h
+++ b/llvm/include/llvm/IR/IRBuilder.h
@@ -2438,12 +2438,14 @@ class IRBuilderBase {
public:
CallInst *CreateCall(FunctionType *FTy, Value *Callee,
ArrayRef<Value *> Args = {}, const Twine &Name = "",
- MDNode *FPMathTag = nullptr) {
+ MDNode *FPMathTag = nullptr,
+ FastMathFlags *uFMF = nullptr) {
CallInst *CI = CallInst::Create(FTy, Callee, Args, DefaultOperandBundles);
if (IsFPConstrained)
setConstrainedFPCallAttr(CI);
- if (isa<FPMathOperator>(CI))
- setFPAttrs(CI, FPMathTag, FMF);
+ if (isa<FPMathOperator>(CI)) {
+ setFPAttrs(CI, FPMathTag, uFMF ? (FMF | *uFMF) : FMF);
+ }
return Insert(CI, Name);
}
@@ -2459,9 +2461,10 @@ class IRBuilderBase {
}
CallInst *CreateCall(FunctionCallee Callee, ArrayRef<Value *> Args = {},
- const Twine &Name = "", MDNode *FPMathTag = nullptr) {
+ const Twine &Name = "", MDNode *FPMathTag = nullptr,
+ FastMathFlags *uFMF = nullptr) {
return CreateCall(Callee.getFunctionType(), Callee.getCallee(), Args, Name,
- FPMathTag);
+ FPMathTag, uFMF);
}
CallInst *CreateCall(FunctionCallee Callee, ArrayRef<Value *> Args,
@@ -2474,7 +2477,8 @@ class IRBuilderBase {
CallInst *CreateConstrainedFPCall(
Function *Callee, ArrayRef<Value *> Args, const Twine &Name = "",
std::optional<RoundingMode> Rounding = std::nullopt,
- std::optional<fp::ExceptionBehavior> Except = std::nullopt);
+ std::optional<fp::ExceptionBehavior> Except = std::nullopt,
+ FastMathFlags *FMF = nullptr);
Value *CreateSelect(Value *C, Value *True, Value *False,
const Twine &Name = "", Instruction *MDFrom = nullptr);
diff --git a/llvm/lib/IR/IRBuilder.cpp b/llvm/lib/IR/IRBuilder.cpp
index f340f7aafdc76f..5feaf956b45a97 100644
--- a/llvm/lib/IR/IRBuilder.cpp
+++ b/llvm/lib/IR/IRBuilder.cpp
@@ -1031,7 +1031,7 @@ CallInst *IRBuilderBase::CreateConstrainedFPCmp(
CallInst *IRBuilderBase::CreateConstrainedFPCall(
Function *Callee, ArrayRef<Value *> Args, const Twine &Name,
std::optional<RoundingMode> Rounding,
- std::optional<fp::ExceptionBehavior> Except) {
+ std::optional<fp::ExceptionBehavior> Except, FastMathFlags *FMF) {
llvm::SmallVector<Value *, 6> UseArgs;
append_range(UseArgs, Args);
@@ -1040,7 +1040,7 @@ CallInst *IRBuilderBase::CreateConstrainedFPCall(
UseArgs.push_back(getConstrainedFPRounding(Rounding));
UseArgs.push_back(getConstrainedFPExcept(Except));
- CallInst *C = CreateCall(Callee, UseArgs, Name);
+ CallInst *C = CreateCall(Callee, UseArgs, Name, nullptr, FMF);
setConstrainedFPCallAttr(C);
return C;
}
>From b6333d0e399a0c7991b597ca00635dbb6ac887aa Mon Sep 17 00:00:00 2001
From: YunQiang Su <yunqiang at isrc.iscas.ac.cn>
Date: Tue, 22 Oct 2024 09:56:06 +0800
Subject: [PATCH 2/2] fix testcases
---
clang/test/CodeGen/fmaxnum_fminnum_use_nsz.c | 36 ++++++++++++---
clang/test/Headers/__clang_hip_math.hip | 48 ++++++++++----------
2 files changed, 54 insertions(+), 30 deletions(-)
diff --git a/clang/test/CodeGen/fmaxnum_fminnum_use_nsz.c b/clang/test/CodeGen/fmaxnum_fminnum_use_nsz.c
index 9798baf0432fea..a708d5cd3ac9e0 100644
--- a/clang/test/CodeGen/fmaxnum_fminnum_use_nsz.c
+++ b/clang/test/CodeGen/fmaxnum_fminnum_use_nsz.c
@@ -8,26 +8,50 @@ double fmax (double, double);
long double fmaxl (long double, long double);
// CHECK: call nsz float @llvm.minnum.f32
-float fmin1(float a, float b) {
+float fmin32(float a, float b) {
return fminf(a, b);
}
+// CHECK: call nsz float @llvm.minnum.f32
+float fmin32b(float a, float b) {
+ return __builtin_fminf(a, b);
+}
// CHECK: call nsz double @llvm.minnum.f64
-float fmin2(double a, double b) {
+float fmin64(double a, double b) {
return fmin(a, b);
}
+// CHECK: call nsz double @llvm.minnum.f64
+float fmin64b(double a, double b) {
+ return __builtin_fmin(a, b);
+}
// CHECK: call nsz x86_fp80 @llvm.minnum.f80
-float fmin3(long double a, long double b) {
+float fmin80(long double a, long double b) {
return fminl(a, b);
}
+// CHECK: call nsz x86_fp80 @llvm.minnum.f80
+float fmin80b(long double a, long double b) {
+ return __builtin_fminl(a, b);
+}
// CHECK: call nsz float @llvm.maxnum.f32
-float fmax1(float a, float b) {
+float fmax32(float a, float b) {
return fmaxf(a, b);
}
+// CHECK: call nsz float @llvm.maxnum.f32
+float fmax32b(float a, float b) {
+ return __builtin_fmaxf(a, b);
+}
// CHECK: call nsz double @llvm.maxnum.f64
-float fmax2(double a, double b) {
+float fmax64(double a, double b) {
return fmax(a, b);
}
+// CHECK: call nsz double @llvm.maxnum.f64
+float fmax64b(double a, double b) {
+ return __builtin_fmax(a, b);
+}
// CHECK: call nsz x86_fp80 @llvm.maxnum.f80
-float fmax3(long double a, long double b) {
+float fmax80(long double a, long double b) {
return fmaxl(a, b);
}
+// CHECK: call nsz x86_fp80 @llvm.maxnum.f80
+float fmax80b(long double a, long double b) {
+ return __builtin_fmaxl(a, b);
+}
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index e4254d1e64bec9..b37f11e16fbd42 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -1248,17 +1248,17 @@ extern "C" __device__ double test_fma_rn(double x, double y, double z) {
// DEFAULT-LABEL: @test_fmaxf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_fmaxf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_fmaxf(
// APPROX-NEXT: entry:
-// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_fmaxf(float x, float y) {
@@ -1267,17 +1267,17 @@ extern "C" __device__ float test_fmaxf(float x, float y) {
// DEFAULT-LABEL: @test_fmax(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_fmax(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_fmax(
// APPROX-NEXT: entry:
-// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_fmax(double x, double y) {
@@ -1286,17 +1286,17 @@ extern "C" __device__ double test_fmax(double x, double y) {
// DEFAULT-LABEL: @test_fminf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_fminf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_fminf(
// APPROX-NEXT: entry:
-// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_fminf(float x, float y) {
@@ -1305,17 +1305,17 @@ extern "C" __device__ float test_fminf(float x, float y) {
// DEFAULT-LABEL: @test_fmin(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_fmin(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_fmin(
// APPROX-NEXT: entry:
-// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_fmin(double x, double y) {
@@ -5114,17 +5114,17 @@ extern "C" __device__ double test__fma_rn(double x, double y, double z) {
// DEFAULT-LABEL: @test_float_min(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_float_min(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_float_min(
// APPROX-NEXT: entry:
-// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_float_min(float x, float y) {
@@ -5133,17 +5133,17 @@ extern "C" __device__ float test_float_min(float x, float y) {
// DEFAULT-LABEL: @test_float_max(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_float_max(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_float_max(
// APPROX-NEXT: entry:
-// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_float_max(float x, float y) {
@@ -5152,17 +5152,17 @@ extern "C" __device__ float test_float_max(float x, float y) {
// DEFAULT-LABEL: @test_double_min(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_double_min(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_double_min(
// APPROX-NEXT: entry:
-// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_double_min(double x, double y) {
@@ -5171,17 +5171,17 @@ extern "C" __device__ double test_double_min(double x, double y) {
// DEFAULT-LABEL: @test_double_max(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_double_max(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_double_max(
// APPROX-NEXT: entry:
-// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
// APPROX-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_double_max(double x, double y) {
More information about the cfe-commits
mailing list