[llvm] [Support][KnownFPClass] Generalize scaling refinement to an exponent range (PR #215094)
Max Graey via llvm-commits
llvm-commits at lists.llvm.org
Mon Aug 10 04:42:56 PDT 2026
https://github.com/MaxGraey updated https://github.com/llvm/llvm-project/pull/215094
>From 4674b67901f344a96b7598832dfdecef587e2970 Mon Sep 17 00:00:00 2001
From: MaxGraey <maxgraey at gmail.com>
Date: Sun, 9 Aug 2026 16:16:37 +0300
Subject: [PATCH 1/5] Generalize scaling refinement to an exponent range
---
llvm/include/llvm/Support/KnownFPClass.h | 14 +-
llvm/lib/Analysis/ValueTracking.cpp | 36 +-
.../CodeGen/GlobalISel/GISelValueTracking.cpp | 24 +-
llvm/lib/Support/KnownFPClass.cpp | 155 ++++++---
.../InstCombineSimplifyDemanded.cpp | 2 +-
.../Transforms/Attributor/nofpclass-fdiv.ll | 305 ++++++++++++++++-
.../Transforms/Attributor/nofpclass-fmul.ll | 188 ++++++++++-
.../Transforms/Attributor/nofpclass-ldexp.ll | 317 ++++++++++++------
llvm/test/Transforms/Attributor/nofpclass.ll | 74 +++-
.../Transforms/InstCombine/fptoui-of-fdiv.ll | 5 +-
.../CodeGen/GlobalISel/KnownFPClassTest.cpp | 23 ++
11 files changed, 925 insertions(+), 218 deletions(-)
diff --git a/llvm/include/llvm/Support/KnownFPClass.h b/llvm/include/llvm/Support/KnownFPClass.h
index 956141b9449a2..ed4af9b18af29 100644
--- a/llvm/include/llvm/Support/KnownFPClass.h
+++ b/llvm/include/llvm/Support/KnownFPClass.h
@@ -248,7 +248,7 @@ struct KnownFPClass {
/// Report known values for fadd x, x
LLVM_ABI static KnownFPClass
- fadd_self(const KnownFPClass &Src,
+ fadd_self(const KnownFPClass &Src, const fltSemantics &Sem,
DenormalMode Mode = DenormalMode::getDynamic());
/// Report known values for fsub
@@ -281,6 +281,10 @@ struct KnownFPClass {
fdiv(const KnownFPClass &LHS, const KnownFPClass &RHS,
DenormalMode Mode = DenormalMode::getDynamic());
+ LLVM_ABI static KnownFPClass
+ fdiv(const KnownFPClass &LHS, const APFloat &RHS,
+ DenormalMode Mode = DenormalMode::getDynamic());
+
/// Report known values for fdiv x, x
LLVM_ABI static KnownFPClass
fdiv_self(const KnownFPClass &Src,
@@ -405,6 +409,14 @@ struct KnownFPClass {
knownNot(fcPositive);
}
+ /// Refine the known classes of Src scaled by a factor with magnitude in
+ /// [2^LoExp, 2^HiExp], as ldexp and a constant multiply or divide do. A
+ /// scale that cannot grow rules out infinities and normals, one that cannot
+ /// shrink rules out subnormals and zeroes. NegativeScale flips the sign.
+ LLVM_ABI void propagateExpRange(const KnownFPClass &Src, int LoExp, int HiExp,
+ bool NegativeScale, const fltSemantics &Sem,
+ DenormalMode Mode);
+
/// Propagate knowledge from a source value that could be a denormal or
/// zero. We have to be conservative since output flushing is not guaranteed,
/// so known-never-zero may not hold.
diff --git a/llvm/lib/Analysis/ValueTracking.cpp b/llvm/lib/Analysis/ValueTracking.cpp
index 44a1240f5635a..66f587f950f6b 100644
--- a/llvm/lib/Analysis/ValueTracking.cpp
+++ b/llvm/lib/Analysis/ValueTracking.cpp
@@ -5592,18 +5592,17 @@ void computeKnownFPClass(const Value *V, const APInt &DemandedElts,
break;
}
case Intrinsic::ldexp: {
+ // Ruling out a zero result also reads the source subnormal classes.
+ FPClassTest InterestedSrcs = InterestedClasses;
+ if (InterestedClasses & fcZero)
+ InterestedSrcs |= fcSubnormal;
+
KnownFPClass KnownSrc;
- computeKnownFPClass(II->getArgOperand(0), DemandedElts, InterestedClasses,
+ computeKnownFPClass(II->getArgOperand(0), DemandedElts, InterestedSrcs,
KnownSrc, Q, Depth + 1);
- // Can refine inf/zero handling based on the exponent operand.
- const FPClassTest ExpInfoMask = fcZero | fcSubnormal | fcInf;
- const Value *ExpArg = II->getArgOperand(1);
- ConstantRange ExpKnownRange =
- ((KnownSrc.KnownFPClasses & ExpInfoMask) != fcNone)
- ? computeConstantRange(ExpArg, /*ForSigned=*/true, Q, Depth + 1)
- : ConstantRange::getFull(
- ExpArg->getType()->getScalarSizeInBits());
+ ConstantRange ExpKnownRange = computeConstantRange(
+ II->getArgOperand(1), /*ForSigned=*/true, Q, Depth + 1);
const fltSemantics &Flt =
II->getType()->getScalarType()->getFltSemantics();
@@ -5811,7 +5810,7 @@ void computeKnownFPClass(const Value *V, const APInt &DemandedElts,
F ? F->getDenormalMode(FltSem) : DenormalMode::getDynamic();
if (Self && Opc == Instruction::FAdd) {
- Known = KnownFPClass::fadd_self(KnownLHS, Mode);
+ Known = KnownFPClass::fadd_self(KnownLHS, FltSem, Mode);
} else {
// RHS is canonically cheaper to compute. Skip inspecting the LHS if
// there's no point.
@@ -5852,6 +5851,7 @@ void computeKnownFPClass(const Value *V, const APInt &DemandedElts,
const APFloat *CRHS;
if (match(RHS, m_APFloat(CRHS))) {
+ KnownRHS = KnownFPClass(*CRHS);
computeKnownFPClass(LHS, DemandedElts, fcAllFlags, KnownLHS, Q,
Depth + 1);
Known = KnownFPClass::fmul(KnownLHS, *CRHS, Mode);
@@ -5913,17 +5913,22 @@ void computeKnownFPClass(const Value *V, const APInt &DemandedElts,
break;
}
+ // A constant divisor bounds the scale factor, which refines both signs.
+ const APFloat *CRHS = nullptr;
+ if (Opc == Instruction::FDiv)
+ match(Op->getOperand(1), m_APFloat(CRHS));
+
const bool WantNegative = (InterestedClasses & fcNegative) != fcNone;
- const bool WantPositive =
- Opc == Instruction::FRem && (InterestedClasses & fcPositive) != fcNone;
+ const bool WantPositive = (InterestedClasses & fcPositive) != fcNone &&
+ (Opc == Instruction::FRem || CRHS);
if (!WantNan && !WantNegative && !WantPositive)
break;
KnownFPClass KnownLHS, KnownRHS;
computeKnownFPClass(Op->getOperand(1), DemandedElts,
- fcNan | fcInf | fcZero | fcNegative, KnownRHS, Q,
- Depth + 1);
+ fcNan | fcInf | fcZero | fcSubnormal | fcNegative,
+ KnownRHS, Q, Depth + 1);
bool KnowSomethingUseful = KnownRHS.isKnownNeverNaN() ||
KnownRHS.isKnownNever(fcNegative) ||
@@ -5941,7 +5946,8 @@ void computeKnownFPClass(const Value *V, const APInt &DemandedElts,
if (Op->getOpcode() == Instruction::FDiv) {
DenormalMode Mode =
F ? F->getDenormalMode(FltSem) : DenormalMode::getDynamic();
- Known = KnownFPClass::fdiv(KnownLHS, KnownRHS, Mode);
+ Known = CRHS ? KnownFPClass::fdiv(KnownLHS, *CRHS, Mode)
+ : KnownFPClass::fdiv(KnownLHS, KnownRHS, Mode);
} else {
// Inf REM x and x REM 0 produce NaN.
if (KnownLHS.isKnownNeverNaN() && KnownRHS.isKnownNeverNaN() &&
diff --git a/llvm/lib/CodeGen/GlobalISel/GISelValueTracking.cpp b/llvm/lib/CodeGen/GlobalISel/GISelValueTracking.cpp
index 139d70f28a4d9..a7491a9a0004c 100644
--- a/llvm/lib/CodeGen/GlobalISel/GISelValueTracking.cpp
+++ b/llvm/lib/CodeGen/GlobalISel/GISelValueTracking.cpp
@@ -1683,8 +1683,8 @@ void GISelValueTracking::computeKnownFPClass(Register R,
break;
}
- DenormalMode Mode =
- MF->getDenormalMode(getFltSemanticForLLT(DstTy.getScalarType()));
+ const fltSemantics &Flt = getFltSemanticForLLT(DstTy.getScalarType());
+ DenormalMode Mode = MF->getDenormalMode(Flt);
FPClassTest InterestedSrcs = InterestedClasses;
if (WantNegative)
@@ -1697,7 +1697,7 @@ void GISelValueTracking::computeKnownFPClass(Register R,
KnownFPClass KnownSelf;
computeKnownFPClass(LHS, DemandedElts, InterestedSrcs, KnownSelf,
Depth + 1);
- Known = KnownFPClass::fadd_self(KnownSelf, Mode);
+ Known = KnownFPClass::fadd_self(KnownSelf, Flt, Mode);
break;
}
@@ -1795,17 +1795,26 @@ void GISelValueTracking::computeKnownFPClass(Register R,
break;
}
+ // A constant divisor bounds the scale factor, which refines both signs.
+ std::optional<APFloat> CRHS;
+ if (Opcode == TargetOpcode::G_FDIV) {
+ auto RHSCst = GFConstant::getConstant(RHS, MRI);
+ if (RHSCst && RHSCst->getKind() == GFConstant::GFConstantKind::Scalar)
+ CRHS = RHSCst->getScalarValue();
+ }
+
const bool WantNan = (InterestedClasses & fcNan) != fcNone;
const bool WantNegative = (InterestedClasses & fcNegative) != fcNone;
- const bool WantPositive = Opcode == TargetOpcode::G_FREM &&
- (InterestedClasses & fcPositive) != fcNone;
+ const bool WantPositive = (InterestedClasses & fcPositive) != fcNone &&
+ (Opcode == TargetOpcode::G_FREM || CRHS);
if (!WantNan && !WantNegative && !WantPositive) {
break;
}
KnownFPClass KnownLHS, KnownRHS;
- computeKnownFPClass(RHS, DemandedElts, fcNan | fcInf | fcZero | fcNegative,
+ computeKnownFPClass(RHS, DemandedElts,
+ fcNan | fcInf | fcZero | fcSubnormal | fcNegative,
KnownRHS, Depth + 1);
bool KnowSomethingUseful = KnownRHS.isKnownNeverNaN() ||
@@ -1817,7 +1826,8 @@ void GISelValueTracking::computeKnownFPClass(Register R,
}
if (Opcode == TargetOpcode::G_FDIV) {
- Known = KnownFPClass::fdiv(KnownLHS, KnownRHS, Mode);
+ Known = CRHS ? KnownFPClass::fdiv(KnownLHS, *CRHS, Mode)
+ : KnownFPClass::fdiv(KnownLHS, KnownRHS, Mode);
} else {
// Inf REM x and x REM 0 produce NaN.
if (KnownLHS.isKnownNeverNaN() && KnownRHS.isKnownNeverNaN() &&
diff --git a/llvm/lib/Support/KnownFPClass.cpp b/llvm/lib/Support/KnownFPClass.cpp
index eccd83451a05f..f9522fadfcf02 100644
--- a/llvm/lib/Support/KnownFPClass.cpp
+++ b/llvm/lib/Support/KnownFPClass.cpp
@@ -13,8 +13,10 @@
#include "llvm/Support/KnownFPClass.h"
#include "llvm/ADT/APFloat.h"
+#include "llvm/ADT/APInt.h"
#include "llvm/Support/ErrorHandling.h"
#include "llvm/Support/KnownBits.h"
+#include <climits>
using namespace llvm;
@@ -65,6 +67,62 @@ bool KnownFPClass::isKnownNeverLogicalPosZero(DenormalMode Mode) const {
llvm_unreachable("covered switch over denormal mode");
}
+void KnownFPClass::propagateExpRange(const KnownFPClass &Src, int LoExp,
+ int HiExp, bool NegativeScale,
+ const fltSemantics &Sem,
+ DenormalMode Mode) {
+ FPClassTest Result = fcNone;
+ FPClassTest KnownFPClasses = Src.KnownFPClasses;
+
+ if (!Src.isKnownNeverLogicalPosZero(Mode))
+ KnownFPClasses |= fcPosZero;
+ if (!Src.isKnownNeverLogicalNegZero(Mode))
+ KnownFPClasses |= fcNegZero;
+
+ if (NegativeScale)
+ KnownFPClasses = llvm::fneg(KnownFPClasses);
+
+ // A scale that cannot grow keeps |result| <= |source|.
+ if (HiExp <= 0)
+ Result |= ~KnownFPClasses & (fcInf | fcNormal);
+
+ if (LoExp >= 0) {
+ // 2^MantissaBits lifts even the smallest value 2^(emin - MantissaBits) into
+ // the normal range. Double-double reports no precision and has no such
+ // threshold, only its source can rule a subnormal out.
+ const unsigned Precision = APFloat::semanticsPrecision(Sem);
+ const bool LiftsOutOfSubnormals =
+ Precision != 0 && LoExp >= static_cast<int>(Precision) - 1;
+
+ Result |=
+ LiftsOutOfSubnormals ? fcSubnormal : ~KnownFPClasses & fcSubnormal;
+
+ // A subnormal result can still be flushed onto a zero
+ //
+ // ieee no flush
+ // preservesign +sub => +0, -sub => -0
+ // positivezero +sub => +0, -sub => +0
+ // dynamic +sub => +0, -sub => +/-0
+
+ FPClassTest FlushToPos = fcNone;
+ FPClassTest FlushToNeg = fcNone;
+
+ if (Mode.Output != DenormalMode::IEEE) {
+ const bool IsKeepSign = Mode.Output == DenormalMode::PreserveSign;
+ const bool IsPosZero = Mode.Output == DenormalMode::PositiveZero;
+ FlushToPos = IsKeepSign ? fcPosSubnormal : fcSubnormal;
+ FlushToNeg = IsPosZero ? fcNone : fcNegSubnormal;
+ }
+
+ if ((Result & FlushToPos) == FlushToPos)
+ Result |= ~KnownFPClasses & fcPosZero;
+ if ((Result & FlushToNeg) == FlushToNeg)
+ Result |= ~KnownFPClasses & fcNegZero;
+ }
+
+ knownNot(Result);
+}
+
void KnownFPClass::propagateDenormal(const KnownFPClass &Src,
DenormalMode Mode) {
KnownFPClasses = Src.KnownFPClasses;
@@ -334,18 +392,17 @@ KnownFPClass KnownFPClass::fadd(const KnownFPClass &KnownLHS,
}
KnownFPClass KnownFPClass::fadd_self(const KnownFPClass &KnownSrc,
+ const fltSemantics &Sem,
DenormalMode Mode) {
KnownFPClass Known = fadd(KnownSrc, KnownSrc, Mode);
- // Doubling 0 will give the same 0.
- if (KnownSrc.isKnownNeverLogicalPosZero(Mode) &&
- (Mode.Output == DenormalMode::IEEE ||
- (Mode.Output == DenormalMode::PreserveSign &&
- KnownSrc.isKnownNeverPosSubnormal()) ||
- (Mode.Output == DenormalMode::PositiveZero &&
- KnownSrc.isKnownNeverSubnormal())))
- Known.knownNot(fcPosZero);
+ // The scale can only refine the subnormal and zero classes here.
+ if ((KnownSrc.KnownFPClasses & (fcZero | fcSubnormal)) ==
+ (fcZero | fcSubnormal))
+ return Known;
+ // x + x scales by 2, down to 1 when one read flushes and leaves x + 0 = x.
+ Known.propagateExpRange(KnownSrc, 0, 1, /*NegativeScale=*/false, Sem, Mode);
return Known;
}
@@ -387,32 +444,20 @@ KnownFPClass KnownFPClass::fmul(const KnownFPClass &KnownLHS,
return Known;
}
-// TODO: This generalizes to known ranges
KnownFPClass KnownFPClass::fmul(const KnownFPClass &KnownLHS,
- const APFloat &CRHS, DenormalMode Mode) {
- // Match denormal scaling pattern, similar to the case in ldexp. If the
- // constant's exponent is sufficiently large, the result cannot be subnormal.
-
- const fltSemantics &Flt = CRHS.getSemantics();
- unsigned Precision = APFloat::semanticsPrecision(Flt);
- const int MantissaBits = Precision - 1;
+ const APFloat &ConstRHS, DenormalMode Mode) {
+ auto Known = KnownFPClass::fmul(KnownLHS, KnownFPClass(ConstRHS), Mode);
- int MinKnownExponent = ilogb(CRHS);
- bool CannotBeSubnormal = (MinKnownExponent >= MantissaBits);
-
- KnownFPClass Known = KnownFPClass::fmul(KnownLHS, KnownFPClass(CRHS), Mode);
- if (CannotBeSubnormal)
- Known.knownNot(fcSubnormal);
+ // ilogb needs a finite nonzero magnitude.
+ if (!ConstRHS.isFiniteNonZero())
+ return Known;
- // Multiply of values <= 1 cannot introduce overflow.
- if (KnownLHS.isKnownNever(fcInf)) {
- if (MinKnownExponent < 0)
- Known.knownNot(fcInf);
- else if (MinKnownExponent == 0 && CRHS.compareAbsoluteValue(APFloat::getOne(
- Flt)) == APFloat::cmpEqual)
- Known.knownNot(fcInf);
- }
+ const int Exp = ilogb(ConstRHS);
+ const bool IsPow2 = ConstRHS.getExactLog2Abs() != INT_MIN;
+ // |C| in [2^Exp, 2^(Exp+1)), exactly 2^Exp for a power of two.
+ Known.propagateExpRange(KnownLHS, Exp, IsPow2 ? Exp : Exp + 1,
+ ConstRHS.isNegative(), ConstRHS.getSemantics(), Mode);
return Known;
}
@@ -445,6 +490,23 @@ KnownFPClass KnownFPClass::fdiv(const KnownFPClass &KnownLHS,
return Known;
}
+KnownFPClass KnownFPClass::fdiv(const KnownFPClass &KnownLHS,
+ const APFloat &ConstRHS, DenormalMode Mode) {
+ auto Known = KnownFPClass::fdiv(KnownLHS, KnownFPClass(ConstRHS), Mode);
+
+ // ilogb needs a finite nonzero magnitude.
+ if (!ConstRHS.isFiniteNonZero())
+ return Known;
+
+ const int Exp = ilogb(ConstRHS);
+ const bool IsPow2 = ConstRHS.getExactLog2Abs() != INT_MIN;
+
+ // |1 / C| in (2^(-Exp-1), 2^-Exp], exactly 2^-Exp for a power of two.
+ Known.propagateExpRange(KnownLHS, IsPow2 ? -Exp : -Exp - 1, -Exp,
+ ConstRHS.isNegative(), ConstRHS.getSemantics(), Mode);
+ return Known;
+}
+
KnownFPClass KnownFPClass::fdiv_self(const KnownFPClass &KnownSrc,
DenormalMode Mode) {
// X / X is always exactly 1.0 or a NaN.
@@ -803,32 +865,23 @@ KnownFPClass KnownFPClass::ldexp(const KnownFPClass &KnownSrc,
else if (KnownSrc.cannotBeOrderedGreaterThanZero())
Known.knownNot(OrderedGreaterThanZeroMask);
- unsigned Precision = APFloat::semanticsPrecision(Flt);
- const int MantissaBits = Precision - 1;
- if (ConstantRangeExpMin.sge(MantissaBits))
- Known.knownNot(fcSubnormal);
-
if (ConstantRangeExpMin.isZero() && ConstantRangeExpMax.isZero()) {
// ldexp(x, 0) -> x, so propagate everything.
Known.propagateCanonicalizingSrc(KnownSrc, Mode);
- } else if (ConstantRangeExpMax.isNonPositive()) {
- // If we know the power is <= 0, can't introduce inf
- if (KnownSrc.isKnownNeverPosInfinity())
- Known.knownNot(fcPosInf);
- if (KnownSrc.isKnownNeverNegInfinity())
- Known.knownNot(fcNegInf);
- } else if (ConstantRangeExpMin.isNonNegative()) {
- // If we know the power is >= 0, can't introduce subnormal or zero
- if (KnownSrc.isKnownNeverPosSubnormal())
- Known.knownNot(fcPosSubnormal);
- if (KnownSrc.isKnownNeverNegSubnormal())
- Known.knownNot(fcNegSubnormal);
- if (KnownSrc.isKnownNeverLogicalPosZero(Mode))
- Known.knownNot(fcPosZero);
- if (KnownSrc.isKnownNeverLogicalNegZero(Mode))
- Known.knownNot(fcNegZero);
+ return Known;
}
+ // The exponent operand can be any integer width. Truncating it would map a
+ // huge exponent onto a small one and invert the tests below.
+ auto SaturatedIntExp = [](const APInt &Exp) {
+ if (Exp.isSignedIntN(sizeof(int) * CHAR_BIT))
+ return static_cast<int>(Exp.getSExtValue());
+ return Exp.isNegative() ? INT_MIN : INT_MAX;
+ };
+
+ Known.propagateExpRange(KnownSrc, SaturatedIntExp(ConstantRangeExpMin),
+ SaturatedIntExp(ConstantRangeExpMax),
+ /*NegativeScale=*/false, Flt, Mode);
return Known;
}
diff --git a/llvm/lib/Transforms/InstCombine/InstCombineSimplifyDemanded.cpp b/llvm/lib/Transforms/InstCombine/InstCombineSimplifyDemanded.cpp
index 0a7c0375106f2..4c93d39f5eb31 100644
--- a/llvm/lib/Transforms/InstCombine/InstCombineSimplifyDemanded.cpp
+++ b/llvm/lib/Transforms/InstCombine/InstCombineSimplifyDemanded.cpp
@@ -2488,7 +2488,7 @@ Value *InstCombinerImpl::SimplifyDemandedUseFPClass(Instruction *I,
Depth + 1))
return I;
- Known = KnownFPClass::fadd_self(KnownLHS, Mode);
+ Known = KnownFPClass::fadd_self(KnownLHS, EltTy->getFltSemantics(), Mode);
KnownRHS = KnownLHS;
} else {
FPClassTest SrcDemandedMask = fcFinite;
diff --git a/llvm/test/Transforms/Attributor/nofpclass-fdiv.ll b/llvm/test/Transforms/Attributor/nofpclass-fdiv.ll
index fef871a1e85f7..d30530e66e72c 100644
--- a/llvm/test/Transforms/Attributor/nofpclass-fdiv.ll
+++ b/llvm/test/Transforms/Attributor/nofpclass-fdiv.ll
@@ -870,7 +870,7 @@ define float @ret_known_zero_or_nan_fdiv_known_inf(float nofpclass(inf norm sub)
define float @ret_fdiv_lhs_known_positive_or_nan(float %lhs, float %rhs) {
; CHECK-LABEL: define float @ret_fdiv_lhs_known_positive_or_nan
; CHECK-SAME: (float [[LHS:%.*]], float [[RHS:%.*]]) #[[ATTR4]] {
-; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR6:[0-9]+]]
+; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR7:[0-9]+]]
; CHECK-NEXT: [[MUL:%.*]] = fdiv float [[LHS_FABS]], [[RHS]]
; CHECK-NEXT: ret float [[MUL]]
;
@@ -883,7 +883,7 @@ define float @ret_fdiv_lhs_known_positive_or_nan(float %lhs, float %rhs) {
define float @ret_fdiv_rhs_known_positive_or_nan(float %lhs, float %rhs) {
; CHECK-LABEL: define float @ret_fdiv_rhs_known_positive_or_nan
; CHECK-SAME: (float [[LHS:%.*]], float [[RHS:%.*]]) #[[ATTR4]] {
-; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR6]]
+; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR7]]
; CHECK-NEXT: [[MUL:%.*]] = fdiv float [[LHS]], [[RHS_FABS]]
; CHECK-NEXT: ret float [[MUL]]
;
@@ -896,8 +896,8 @@ define float @ret_fdiv_rhs_known_positive_or_nan(float %lhs, float %rhs) {
define float @ret_fdiv_both_signs_positive_or_nan(float %lhs, float %rhs) {
; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) float @ret_fdiv_both_signs_positive_or_nan
; CHECK-SAME: (float [[LHS:%.*]], float [[RHS:%.*]]) #[[ATTR4]] {
-; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR6]]
-; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR6]]
+; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR7]]
+; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR7]]
; CHECK-NEXT: [[MUL:%.*]] = fdiv float [[LHS_FABS]], [[RHS_FABS]]
; CHECK-NEXT: ret float [[MUL]]
;
@@ -911,8 +911,8 @@ define float @ret_fdiv_both_signs_positive_or_nan(float %lhs, float %rhs) {
define float @ret_fdiv_both_signs_negative_or_nan(float %lhs, float %rhs) {
; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) float @ret_fdiv_both_signs_negative_or_nan
; CHECK-SAME: (float [[LHS:%.*]], float [[RHS:%.*]]) #[[ATTR4]] {
-; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR6]]
-; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR6]]
+; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR7]]
+; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR7]]
; CHECK-NEXT: [[LHS_NEG_FABS:%.*]] = fneg float [[LHS_FABS]]
; CHECK-NEXT: [[RHS_NEG_FABS:%.*]] = fneg float [[RHS_FABS]]
; CHECK-NEXT: [[MUL:%.*]] = fdiv float [[LHS_NEG_FABS]], [[RHS_NEG_FABS]]
@@ -930,8 +930,8 @@ define float @ret_fdiv_both_signs_negative_or_nan(float %lhs, float %rhs) {
define float @ret_fdiv_lhs_negative_rhs_positive(float %lhs, float %rhs) {
; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) float @ret_fdiv_lhs_negative_rhs_positive
; CHECK-SAME: (float [[LHS:%.*]], float [[RHS:%.*]]) #[[ATTR4]] {
-; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR6]]
-; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR6]]
+; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR7]]
+; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR7]]
; CHECK-NEXT: [[LHS_NEG_FABS:%.*]] = fneg float [[LHS_FABS]]
; CHECK-NEXT: [[MUL:%.*]] = fdiv float [[LHS_NEG_FABS]], [[RHS_FABS]]
; CHECK-NEXT: ret float [[MUL]]
@@ -947,8 +947,8 @@ define float @ret_fdiv_lhs_negative_rhs_positive(float %lhs, float %rhs) {
define float @ret_fdiv_rhs_negative_lhs_positive(float %lhs, float %rhs) {
; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) float @ret_fdiv_rhs_negative_lhs_positive
; CHECK-SAME: (float [[LHS:%.*]], float [[RHS:%.*]]) #[[ATTR4]] {
-; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR6]]
-; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR6]]
+; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR7]]
+; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR7]]
; CHECK-NEXT: [[RHS_NEG_FABS:%.*]] = fneg float [[RHS_FABS]]
; CHECK-NEXT: [[MUL:%.*]] = fdiv float [[LHS_FABS]], [[RHS_NEG_FABS]]
; CHECK-NEXT: ret float [[MUL]]
@@ -1056,9 +1056,294 @@ define float @ret_known_inf_or_nan_fdiv_known_inf_or_nan(float nofpclass(norm su
ret float %fdiv
}
+; Dividing by a constant of magnitude >= 1 cannot overflow.
+define float @ret_fdiv_noinf_const_1_5(float nofpclass(inf) %arg0) {
+; CHECK-LABEL: define nofpclass(inf) float @ret_fdiv_noinf_const_1_5
+; CHECK-SAME: (float nofpclass(inf) [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 1.500000e+00
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 1.5
+ ret float %fdiv
+}
+
+; Negative test, 1 / 0.75 is greater than one.
+define float @ret_fdiv_noinf_const_0_75(float nofpclass(inf) %arg0) {
+; CHECK-LABEL: define float @ret_fdiv_noinf_const_0_75
+; CHECK-SAME: (float nofpclass(inf) [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 7.500000e-01
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 0.75
+ ret float %fdiv
+}
+
+define float @ret_fdiv_nopinf_const_3(float nofpclass(pinf) %arg0) {
+; CHECK-LABEL: define nofpclass(pinf) float @ret_fdiv_nopinf_const_3
+; CHECK-SAME: (float nofpclass(pinf) [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 3.000000e+00
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 3.0
+ ret float %fdiv
+}
+
+; Dividing by a constant of magnitude <= 1 cannot underflow.
+define float @ret_fdiv_nozerosub_const_1(float nofpclass(zero sub) %arg0) {
+; CHECK-LABEL: define nofpclass(zero sub) float @ret_fdiv_nozerosub_const_1
+; CHECK-SAME: (float nofpclass(zero sub) [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 1.000000e+00
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 1.0
+ ret float %fdiv
+}
+
+; Negative test, 1 / 1.5 is less than one.
+define float @ret_fdiv_nozerosub_const_1_5(float nofpclass(zero sub) %arg0) {
+; CHECK-LABEL: define float @ret_fdiv_nozerosub_const_1_5
+; CHECK-SAME: (float nofpclass(zero sub) [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 1.500000e+00
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 1.5
+ ret float %fdiv
+}
+
+; A power of two divides exactly. The bound is a single exponent, not a range.
+define float @ret_fdiv_noinf_const_2(float nofpclass(inf) %arg0) {
+; CHECK-LABEL: define nofpclass(inf) float @ret_fdiv_noinf_const_2
+; CHECK-SAME: (float nofpclass(inf) [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 2.000000e+00
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 2.0
+ ret float %fdiv
+}
+
+define float @ret_fdiv_nozerosub_const_0_5(float nofpclass(zero sub) %arg0) {
+; CHECK-LABEL: define nofpclass(zero sub) float @ret_fdiv_nozerosub_const_0_5
+; CHECK-SAME: (float nofpclass(zero sub) [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 5.000000e-01
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 0.5
+ ret float %fdiv
+}
+
+; The exponent comes from the magnitude. A negative power of two keeps the
+; exact bound and only mirrors the sign.
+define float @ret_fdiv_nozero_nopsub_const_neg1(float nofpclass(zero psub) %arg0) {
+; CHECK-LABEL: define nofpclass(zero nsub) float @ret_fdiv_nozero_nopsub_const_neg1
+; CHECK-SAME: (float nofpclass(zero psub) [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], -1.000000e+00
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, -1.0
+ ret float %fdiv
+}
+
+define float @ret_fdiv_nopsub_const_0_75(float nofpclass(psub) %arg0) {
+; CHECK-LABEL: define nofpclass(psub) float @ret_fdiv_nopsub_const_0_75
+; CHECK-SAME: (float nofpclass(psub) [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 7.500000e-01
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 0.75
+ ret float %fdiv
+}
+
+; Negative test, a subnormal source reads as either zero under this mode.
+define float @ret_fdiv_nozero_const_0_75_daz(float nofpclass(zero) %arg0) #1 {
+; CHECK-LABEL: define float @ret_fdiv_nozero_const_0_75_daz
+; CHECK-SAME: (float nofpclass(zero) [[ARG0:%.*]]) #[[ATTR1]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 7.500000e-01
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 0.75
+ ret float %fdiv
+}
+
+; A positive zero input mode reads -sub as +0. Only -0.0 is ruled out.
+define float @ret_fdiv_nozero_const_0_75_dapz(float nofpclass(zero) %arg0) #2 {
+; CHECK-LABEL: define nofpclass(nzero) float @ret_fdiv_nozero_const_0_75_dapz
+; CHECK-SAME: (float nofpclass(zero) [[ARG0:%.*]]) #[[ATTR2]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 7.500000e-01
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 0.75
+ ret float %fdiv
+}
+
+; The same, mirrored by a negative divisor. Only +0.0 is ruled out.
+define float @ret_fdiv_nozero_const_neg0_75_dapz(float nofpclass(zero) %arg0) #2 {
+; CHECK-LABEL: define nofpclass(pzero) float @ret_fdiv_nozero_const_neg0_75_dapz
+; CHECK-SAME: (float nofpclass(zero) [[ARG0:%.*]]) #[[ATTR2]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], -7.500000e-01
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, -0.75
+ ret float %fdiv
+}
+
+; Negative test, a dynamic input mode may read a subnormal as either zero.
+define float @ret_fdiv_nozero_const_0_75_dadyn(float nofpclass(zero) %arg0) #3 {
+; CHECK-LABEL: define float @ret_fdiv_nozero_const_0_75_dadyn
+; CHECK-SAME: (float nofpclass(zero) [[ARG0:%.*]]) #[[ATTR3]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 7.500000e-01
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 0.75
+ ret float %fdiv
+}
+
+; The constant divisor is matched through a vector splat as well.
+define <2 x float> @ret_fdiv_v2f32_nozerosub_splat_0_5(<2 x float> nofpclass(zero sub) %arg0) {
+; CHECK-LABEL: define nofpclass(zero sub) <2 x float> @ret_fdiv_v2f32_nozerosub_splat_0_5
+; CHECK-SAME: (<2 x float> nofpclass(zero sub) [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv <2 x float> [[ARG0]], splat (float 5.000000e-01)
+; CHECK-NEXT: ret <2 x float> [[FDIV]]
+;
+ %fdiv = fdiv <2 x float> %arg0, splat (float 0.5)
+ ret <2 x float> %fdiv
+}
+
+; Scaling up by at least 2^MantissaBits leaves the subnormal range.
+define float @ret_fdiv_exponent_f32_neg23(float %arg0) {
+; CHECK-LABEL: define nofpclass(sub) float @ret_fdiv_exponent_f32_neg23
+; CHECK-SAME: (float [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], f0x34000000
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 0x3E80000000000000
+ ret float %fdiv
+}
+
+; Negative test, 1 / 0x1.8p-23 falls one short of 2^23.
+define float @ret_fdiv_exponent_f32_neg23_nonpow2(float %arg0) {
+; CHECK-LABEL: define float @ret_fdiv_exponent_f32_neg23_nonpow2
+; CHECK-SAME: (float [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], f0x34400000
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 0x3E88000000000000
+ ret float %fdiv
+}
+
+; The threshold follows the format, 2^10 suffices for half.
+define half @ret_fdiv_exponent_f16_neg10(half %arg0) {
+; CHECK-LABEL: define nofpclass(sub) half @ret_fdiv_exponent_f16_neg10
+; CHECK-SAME: (half [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv half [[ARG0]], 9.765620e-04
+; CHECK-NEXT: ret half [[FDIV]]
+;
+ %fdiv = fdiv half %arg0, 0xH1400
+ ret half %fdiv
+}
+
+; Negative test, 1 / 0x1.8p-10 falls one short of 2^10.
+define half @ret_fdiv_exponent_f16_neg10_nonpow2(half %arg0) {
+; CHECK-LABEL: define half @ret_fdiv_exponent_f16_neg10_nonpow2
+; CHECK-SAME: (half [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv half [[ARG0]], 1.464840e-03
+; CHECK-NEXT: ret half [[FDIV]]
+;
+ %fdiv = fdiv half %arg0, 0xH1600
+ ret half %fdiv
+}
+
+; A divisor of magnitude >= 1 cannot grow the result. A normal or an infinity
+; must come from a source of the same class.
+define float @ret_fdiv_nonorminf_const_4(float nofpclass(norm inf) %arg0) {
+; CHECK-LABEL: define nofpclass(inf norm) float @ret_fdiv_nonorminf_const_4
+; CHECK-SAME: (float nofpclass(inf norm) [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 4.000000e+00
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 4.0
+ ret float %fdiv
+}
+
+; Negative test, dividing by 0.75 can lift a subnormal into the normal range.
+define float @ret_fdiv_nonorminf_const_0_75(float nofpclass(norm inf) %arg0) {
+; CHECK-LABEL: define float @ret_fdiv_nonorminf_const_0_75
+; CHECK-SAME: (float nofpclass(inf norm) [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 7.500000e-01
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 0.75
+ ret float %fdiv
+}
+
+; Negative test, double-double reports no precision, so no scale is known to
+; leave the subnormal range.
+define ppc_fp128 @ret_fdiv_ppcf128_const_0_5(ppc_fp128 %arg0) {
+; CHECK-LABEL: define ppc_fp128 @ret_fdiv_ppcf128_const_0_5
+; CHECK-SAME: (ppc_fp128 [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv ppc_fp128 [[ARG0]], 5.000000e-01
+; CHECK-NEXT: ret ppc_fp128 [[FDIV]]
+;
+ %fdiv = fdiv ppc_fp128 %arg0, 0xM3FE0000000000000000000000000000
+ ret ppc_fp128 %fdiv
+}
+
+; Negative test, a subnormal result may still be flushed to zero on output.
+define float @ret_fdiv_nozero_const_0_75_ftz_out(float nofpclass(zero) %arg0) #4 {
+; CHECK-LABEL: define float @ret_fdiv_nozero_const_0_75_ftz_out
+; CHECK-SAME: (float nofpclass(zero) [[ARG0:%.*]]) #[[ATTR5:[0-9]+]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 7.500000e-01
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 0.75
+ ret float %fdiv
+}
+
+; A source that is never subnormal cannot reach a flushed zero either.
+define float @ret_fdiv_nozerosub_const_0_75_ftz_out(float nofpclass(zero sub) %arg0) #4 {
+; CHECK-LABEL: define nofpclass(zero sub) float @ret_fdiv_nozerosub_const_0_75_ftz_out
+; CHECK-SAME: (float nofpclass(zero sub) [[ARG0:%.*]]) #[[ATTR5]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 7.500000e-01
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 0.75
+ ret float %fdiv
+}
+
+; A divisor with no exponent to report is left to the class rules alone.
+define float @ret_fdiv_const_zero(float %arg0) {
+; CHECK-LABEL: define nofpclass(zero sub norm) float @ret_fdiv_const_zero
+; CHECK-SAME: (float [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], 0.000000e+00
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 0.0
+ ret float %fdiv
+}
+
+define float @ret_fdiv_const_inf(float %arg0) {
+; CHECK-LABEL: define float @ret_fdiv_const_inf
+; CHECK-SAME: (float [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], +inf
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 0x7FF0000000000000
+ ret float %fdiv
+}
+
+define float @ret_fdiv_const_nan(float %arg0) {
+; CHECK-LABEL: define float @ret_fdiv_const_nan
+; CHECK-SAME: (float [[ARG0:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[FDIV:%.*]] = fdiv float [[ARG0]], +qnan
+; CHECK-NEXT: ret float [[FDIV]]
+;
+ %fdiv = fdiv float %arg0, 0x7FF8000000000000
+ ret float %fdiv
+}
+
attributes #0 = { denormal_fpenv(ieee|ieee) }
attributes #1 = { denormal_fpenv(ieee|preservesign) }
attributes #2 = { denormal_fpenv(ieee|positivezero) }
attributes #3 = { denormal_fpenv(ieee|dynamic) }
+attributes #4 = { denormal_fpenv(preservesign|ieee) }
;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
; TUNIT: {{.*}}
diff --git a/llvm/test/Transforms/Attributor/nofpclass-fmul.ll b/llvm/test/Transforms/Attributor/nofpclass-fmul.ll
index ba98a804afcb4..c2aa74b3a463d 100644
--- a/llvm/test/Transforms/Attributor/nofpclass-fmul.ll
+++ b/llvm/test/Transforms/Attributor/nofpclass-fmul.ll
@@ -618,7 +618,7 @@ define float @ret_known_zero_or_nan_mul_known_inf(float nofpclass(inf norm sub)
define float @ret_fmul_lhs_known_positive_or_nan(float %lhs, float %rhs) {
; CHECK-LABEL: define float @ret_fmul_lhs_known_positive_or_nan(
; CHECK-SAME: float [[LHS:%.*]], float [[RHS:%.*]]) #[[ATTR0]] {
-; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR2:[0-9]+]]
+; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR5:[0-9]+]]
; CHECK-NEXT: [[MUL:%.*]] = fmul float [[LHS_FABS]], [[RHS]]
; CHECK-NEXT: ret float [[MUL]]
;
@@ -631,7 +631,7 @@ define float @ret_fmul_lhs_known_positive_or_nan(float %lhs, float %rhs) {
define float @ret_fmul_rhs_known_positive_or_nan(float %lhs, float %rhs) {
; CHECK-LABEL: define float @ret_fmul_rhs_known_positive_or_nan(
; CHECK-SAME: float [[LHS:%.*]], float [[RHS:%.*]]) #[[ATTR0]] {
-; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR2]]
+; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR5]]
; CHECK-NEXT: [[MUL:%.*]] = fmul float [[LHS]], [[RHS_FABS]]
; CHECK-NEXT: ret float [[MUL]]
;
@@ -644,8 +644,8 @@ define float @ret_fmul_rhs_known_positive_or_nan(float %lhs, float %rhs) {
define float @ret_fmul_both_signs_positive_or_nan(float %lhs, float %rhs) {
; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) float @ret_fmul_both_signs_positive_or_nan(
; CHECK-SAME: float [[LHS:%.*]], float [[RHS:%.*]]) #[[ATTR0]] {
-; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR2]]
-; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR2]]
+; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR5]]
+; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR5]]
; CHECK-NEXT: [[MUL:%.*]] = fmul float [[LHS_FABS]], [[RHS_FABS]]
; CHECK-NEXT: ret float [[MUL]]
;
@@ -659,8 +659,8 @@ define float @ret_fmul_both_signs_positive_or_nan(float %lhs, float %rhs) {
define float @ret_fmul_both_signs_negative_or_nan(float %lhs, float %rhs) {
; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) float @ret_fmul_both_signs_negative_or_nan(
; CHECK-SAME: float [[LHS:%.*]], float [[RHS:%.*]]) #[[ATTR0]] {
-; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR2]]
-; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR2]]
+; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR5]]
+; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR5]]
; CHECK-NEXT: [[LHS_NEG_FABS:%.*]] = fneg float [[LHS_FABS]]
; CHECK-NEXT: [[RHS_NEG_FABS:%.*]] = fneg float [[RHS_FABS]]
; CHECK-NEXT: [[MUL:%.*]] = fmul float [[LHS_NEG_FABS]], [[RHS_NEG_FABS]]
@@ -678,8 +678,8 @@ define float @ret_fmul_both_signs_negative_or_nan(float %lhs, float %rhs) {
define float @ret_fmul_lhs_negative_rhs_positive(float %lhs, float %rhs) {
; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) float @ret_fmul_lhs_negative_rhs_positive(
; CHECK-SAME: float [[LHS:%.*]], float [[RHS:%.*]]) #[[ATTR0]] {
-; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR2]]
-; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR2]]
+; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR5]]
+; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR5]]
; CHECK-NEXT: [[LHS_NEG_FABS:%.*]] = fneg float [[LHS_FABS]]
; CHECK-NEXT: [[MUL:%.*]] = fmul float [[LHS_NEG_FABS]], [[RHS_FABS]]
; CHECK-NEXT: ret float [[MUL]]
@@ -695,8 +695,8 @@ define float @ret_fmul_lhs_negative_rhs_positive(float %lhs, float %rhs) {
define float @ret_fmul_rhs_negative_lhs_positive(float %lhs, float %rhs) {
; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) float @ret_fmul_rhs_negative_lhs_positive(
; CHECK-SAME: float [[LHS:%.*]], float [[RHS:%.*]]) #[[ATTR0]] {
-; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR2]]
-; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR2]]
+; CHECK-NEXT: [[LHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[LHS]]) #[[ATTR5]]
+; CHECK-NEXT: [[RHS_FABS:%.*]] = call float @llvm.fabs.f32(float [[RHS]]) #[[ATTR5]]
; CHECK-NEXT: [[RHS_NEG_FABS:%.*]] = fneg float [[RHS_FABS]]
; CHECK-NEXT: [[MUL:%.*]] = fmul float [[LHS_FABS]], [[RHS_NEG_FABS]]
; CHECK-NEXT: ret float [[MUL]]
@@ -827,7 +827,7 @@ define float @ret_fmul__not_inf__neghalf(float nofpclass(inf) %x) {
}
define float @ret_fmul__not_pinf__half(float nofpclass(pinf) %x) {
-; CHECK-LABEL: define float @ret_fmul__not_pinf__half(
+; CHECK-LABEL: define nofpclass(pinf) float @ret_fmul__not_pinf__half(
; CHECK-SAME: float nofpclass(pinf) [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], 5.000000e-01
; CHECK-NEXT: ret float [[MUL]]
@@ -837,7 +837,7 @@ define float @ret_fmul__not_pinf__half(float nofpclass(pinf) %x) {
}
define float @ret_fmul__not_ninf__half(float nofpclass(ninf) %x) {
-; CHECK-LABEL: define float @ret_fmul__not_ninf__half(
+; CHECK-LABEL: define nofpclass(ninf) float @ret_fmul__not_ninf__half(
; CHECK-SAME: float nofpclass(ninf) [[X:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], 5.000000e-01
; CHECK-NEXT: ret float [[MUL]]
@@ -914,7 +914,7 @@ define float @ret_fmul__not_inf__neg1.5(float nofpclass(inf) %x) {
define float @ret_fmul__not_inf__fsub_floor_pat(float nofpclass(inf) %x, float %y) {
; CHECK-LABEL: define nofpclass(inf) float @ret_fmul__not_inf__fsub_floor_pat(
; CHECK-SAME: float nofpclass(inf) [[X:%.*]], float [[Y:%.*]]) #[[ATTR0]] {
-; CHECK-NEXT: [[FLOOR_Y:%.*]] = call float @llvm.floor.f32(float [[Y]]) #[[ATTR2]]
+; CHECK-NEXT: [[FLOOR_Y:%.*]] = call float @llvm.floor.f32(float [[Y]]) #[[ATTR5]]
; CHECK-NEXT: [[Y_SUB_FLOOR_Y:%.*]] = fsub float [[Y]], [[FLOOR_Y]]
; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], [[Y_SUB_FLOOR_Y]]
; CHECK-NEXT: ret float [[MUL]]
@@ -928,7 +928,7 @@ define float @ret_fmul__not_inf__fsub_floor_pat(float nofpclass(inf) %x, float %
define float @ret_fmul__not_inf__fsub_floor_pat_commute(float nofpclass(inf) %x, float %y) {
; CHECK-LABEL: define nofpclass(inf) float @ret_fmul__not_inf__fsub_floor_pat_commute(
; CHECK-SAME: float nofpclass(inf) [[X:%.*]], float [[Y:%.*]]) #[[ATTR0]] {
-; CHECK-NEXT: [[FLOOR_Y:%.*]] = call float @llvm.floor.f32(float [[Y]]) #[[ATTR2]]
+; CHECK-NEXT: [[FLOOR_Y:%.*]] = call float @llvm.floor.f32(float [[Y]]) #[[ATTR5]]
; CHECK-NEXT: [[Y_SUB_FLOOR_Y:%.*]] = fsub float [[Y]], [[FLOOR_Y]]
; CHECK-NEXT: [[MUL:%.*]] = fmul float [[Y_SUB_FLOOR_Y]], [[X]]
; CHECK-NEXT: ret float [[MUL]]
@@ -942,7 +942,7 @@ define float @ret_fmul__not_inf__fsub_floor_pat_commute(float nofpclass(inf) %x,
define float @ret_fmul__not_inf__fsub_floor_pat_wrong_floor_val(float nofpclass(inf) %x, float %y, float %z) {
; CHECK-LABEL: define float @ret_fmul__not_inf__fsub_floor_pat_wrong_floor_val(
; CHECK-SAME: float nofpclass(inf) [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]]) #[[ATTR0]] {
-; CHECK-NEXT: [[FLOOR_Y:%.*]] = call float @llvm.floor.f32(float [[Z]]) #[[ATTR2]]
+; CHECK-NEXT: [[FLOOR_Y:%.*]] = call float @llvm.floor.f32(float [[Z]]) #[[ATTR5]]
; CHECK-NEXT: [[Y_SUB_FLOOR_Y:%.*]] = fsub float [[Y]], [[FLOOR_Y]]
; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], [[Y_SUB_FLOOR_Y]]
; CHECK-NEXT: ret float [[MUL]]
@@ -952,3 +952,161 @@ define float @ret_fmul__not_inf__fsub_floor_pat_wrong_floor_val(float nofpclass(
%mul = fmul float %x, %y.sub.floor.y
ret float %mul
}
+
+; The same bound applies when the other operand is a constant.
+define float @ret_fmul__not_inf__fsub_floor_pat_const(float %y) {
+; CHECK-LABEL: define nofpclass(inf) float @ret_fmul__not_inf__fsub_floor_pat_const(
+; CHECK-SAME: float [[Y:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT: [[FLOOR_Y:%.*]] = call float @llvm.floor.f32(float [[Y]]) #[[ATTR5]]
+; CHECK-NEXT: [[Y_SUB_FLOOR_Y:%.*]] = fsub float [[Y]], [[FLOOR_Y]]
+; CHECK-NEXT: [[MUL:%.*]] = fmul float [[Y_SUB_FLOOR_Y]], 3.000000e+00
+; CHECK-NEXT: ret float [[MUL]]
+;
+ %floor.y = call float @llvm.floor.f32(float %y)
+ %y.sub.floor.y = fsub float %y, %floor.y
+ %mul = fmul float %y.sub.floor.y, 3.0
+ ret float %mul
+}
+
+; Multiplying by a constant of magnitude >= 1 cannot underflow.
+define float @ret_fmul_nozerosub_const_1_5(float nofpclass(zero sub) %x) {
+; CHECK-LABEL: define nofpclass(zero sub) float @ret_fmul_nozerosub_const_1_5(
+; CHECK-SAME: float nofpclass(zero sub) [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], 1.500000e+00
+; CHECK-NEXT: ret float [[MUL]]
+;
+ %mul = fmul float %x, 1.5
+ ret float %mul
+}
+
+; Negative test, 0.75 is less than one.
+define float @ret_fmul_nozerosub_const_0_75(float nofpclass(zero sub) %x) {
+; CHECK-LABEL: define float @ret_fmul_nozerosub_const_0_75(
+; CHECK-SAME: float nofpclass(zero sub) [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], 7.500000e-01
+; CHECK-NEXT: ret float [[MUL]]
+;
+ %mul = fmul float %x, 0.75
+ ret float %mul
+}
+
+; The sign of the result is mirrored by a negative factor.
+define float @ret_fmul_nopsub_const_neg3(float nofpclass(psub) %x) {
+; CHECK-LABEL: define nofpclass(nsub) float @ret_fmul_nopsub_const_neg3(
+; CHECK-SAME: float nofpclass(psub) [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], -3.000000e+00
+; CHECK-NEXT: ret float [[MUL]]
+;
+ %mul = fmul float %x, -3.0
+ ret float %mul
+}
+
+; Negative test, a subnormal result may still be flushed to zero on output.
+define float @ret_fmul_nozero_const_3_ftz_out(float nofpclass(zero) %x) denormal_fpenv(preservesign|ieee) {
+; CHECK-LABEL: define float @ret_fmul_nozero_const_3_ftz_out(
+; CHECK-SAME: float nofpclass(zero) [[X:%.*]]) #[[ATTR1:[0-9]+]] {
+; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], 3.000000e+00
+; CHECK-NEXT: ret float [[MUL]]
+;
+ %mul = fmul float %x, 3.0
+ ret float %mul
+}
+
+; A source that is never subnormal cannot reach a flushed zero either.
+define float @ret_fmul_nozerosub_const_3_ftz_out(float nofpclass(zero sub) %x) denormal_fpenv(preservesign|ieee) {
+; CHECK-LABEL: define nofpclass(zero sub) float @ret_fmul_nozerosub_const_3_ftz_out(
+; CHECK-SAME: float nofpclass(zero sub) [[X:%.*]]) #[[ATTR1]] {
+; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], 3.000000e+00
+; CHECK-NEXT: ret float [[MUL]]
+;
+ %mul = fmul float %x, 3.0
+ ret float %mul
+}
+
+; A negative factor mirrors the source onto the result. The same four output
+; modes separate the opposite sign of subnormal.
+define float @ret_fmul_nozero_nonsub_const_neg3_ieee_out(float nofpclass(zero nsub) %x) {
+; CHECK-LABEL: define nofpclass(zero psub) float @ret_fmul_nozero_nonsub_const_neg3_ieee_out(
+; CHECK-SAME: float nofpclass(zero nsub) [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], -3.000000e+00
+; CHECK-NEXT: ret float [[MUL]]
+;
+ %mul = fmul float %x, -3.0
+ ret float %mul
+}
+
+define float @ret_fmul_nozero_nonsub_const_neg3_preservesign_out(float nofpclass(zero nsub) %x) denormal_fpenv(preservesign|ieee) {
+; CHECK-LABEL: define nofpclass(pzero psub) float @ret_fmul_nozero_nonsub_const_neg3_preservesign_out(
+; CHECK-SAME: float nofpclass(zero nsub) [[X:%.*]]) #[[ATTR1]] {
+; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], -3.000000e+00
+; CHECK-NEXT: ret float [[MUL]]
+;
+ %mul = fmul float %x, -3.0
+ ret float %mul
+}
+
+define float @ret_fmul_nozero_nonsub_const_neg3_positivezero_out(float nofpclass(zero nsub) %x) denormal_fpenv(positivezero|ieee) {
+; CHECK-LABEL: define nofpclass(nzero psub) float @ret_fmul_nozero_nonsub_const_neg3_positivezero_out(
+; CHECK-SAME: float nofpclass(zero nsub) [[X:%.*]]) #[[ATTR2:[0-9]+]] {
+; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], -3.000000e+00
+; CHECK-NEXT: ret float [[MUL]]
+;
+ %mul = fmul float %x, -3.0
+ ret float %mul
+}
+
+define float @ret_fmul_nozero_nonsub_const_neg3_dynamic_out(float nofpclass(zero nsub) %x) denormal_fpenv(dynamic|ieee) {
+; CHECK-LABEL: define nofpclass(psub) float @ret_fmul_nozero_nonsub_const_neg3_dynamic_out(
+; CHECK-SAME: float nofpclass(zero nsub) [[X:%.*]]) #[[ATTR3:[0-9]+]] {
+; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], -3.000000e+00
+; CHECK-NEXT: ret float [[MUL]]
+;
+ %mul = fmul float %x, -3.0
+ ret float %mul
+}
+
+; A factor of magnitude <= 1 cannot grow the result. A normal or an infinity
+; must come from a source of the same class.
+define float @ret_fmul_nonorminf_const_0_5(float nofpclass(norm inf) %x) {
+; CHECK-LABEL: define nofpclass(inf norm) float @ret_fmul_nonorminf_const_0_5(
+; CHECK-SAME: float nofpclass(inf norm) [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], 5.000000e-01
+; CHECK-NEXT: ret float [[MUL]]
+;
+ %mul = fmul float %x, 0.5
+ ret float %mul
+}
+
+; The same, mirrored by a negative factor.
+define float @ret_fmul_nopnorm_nopinf_const_neg0_5(float nofpclass(pnorm pinf) %x) {
+; CHECK-LABEL: define nofpclass(ninf nnorm) float @ret_fmul_nopnorm_nopinf_const_neg0_5(
+; CHECK-SAME: float nofpclass(pinf pnorm) [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], -5.000000e-01
+; CHECK-NEXT: ret float [[MUL]]
+;
+ %mul = fmul float %x, -0.5
+ ret float %mul
+}
+
+; Negative test, 1.5 can lift a subnormal into the normal range.
+define float @ret_fmul_nonorminf_const_1_5(float nofpclass(norm inf) %x) {
+; CHECK-LABEL: define float @ret_fmul_nonorminf_const_1_5(
+; CHECK-SAME: float nofpclass(inf norm) [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT: [[MUL:%.*]] = fmul float [[X]], 1.500000e+00
+; CHECK-NEXT: ret float [[MUL]]
+;
+ %mul = fmul float %x, 1.5
+ ret float %mul
+}
+
+; Negative test, double-double reports no precision, so no scale is known to
+; leave the subnormal range.
+define ppc_fp128 @ret_fmul_ppcf128_const_2(ppc_fp128 %x) {
+; CHECK-LABEL: define ppc_fp128 @ret_fmul_ppcf128_const_2(
+; CHECK-SAME: ppc_fp128 [[X:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT: [[MUL:%.*]] = fmul ppc_fp128 [[X]], 2.000000e+00
+; CHECK-NEXT: ret ppc_fp128 [[MUL]]
+;
+ %mul = fmul ppc_fp128 %x, 0xM40000000000000000000000000000000
+ ret ppc_fp128 %mul
+}
diff --git a/llvm/test/Transforms/Attributor/nofpclass-ldexp.ll b/llvm/test/Transforms/Attributor/nofpclass-ldexp.ll
index 630e510a12d72..087e10f3bac7a 100644
--- a/llvm/test/Transforms/Attributor/nofpclass-ldexp.ll
+++ b/llvm/test/Transforms/Attributor/nofpclass-ldexp.ll
@@ -7,12 +7,13 @@ declare <2 x float> @llvm.ldexp.v2f32.v2i32(<2 x float>, <2 x i32>)
declare double @llvm.ldexp.f64.i32(double, i32)
declare half @llvm.ldexp.f16.i32(half, i32)
declare bfloat @llvm.ldexp.bf16.i32(bfloat, i32)
+declare ppc_fp128 @llvm.ldexp.ppcf128.i32(ppc_fp128, i32)
declare i32 @llvm.smin.i32(i32, i32)
define float @ret_ldexp_f32(float %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32
; CHECK-SAME: (float [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1:[0-9]+]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 [[ARG1]]) #[[ATTR10:[0-9]+]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 [[ARG1]]) #[[ATTR12:[0-9]+]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -22,7 +23,7 @@ define float @ret_ldexp_f32(float %arg0, i32 %arg1) #0 {
define float @ret_ldexp_f32_0(float %arg0) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_0
; CHECK-SAME: (float [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -32,7 +33,7 @@ define float @ret_ldexp_f32_0(float %arg0) #0 {
define float @ret_ldexp_f32_0_nopsub(float nofpclass(psub) %arg0) #0 {
; CHECK-LABEL: define nofpclass(psub) float @ret_ldexp_f32_0_nopsub
; CHECK-SAME: (float nofpclass(psub) [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(psub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(psub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -42,7 +43,7 @@ define float @ret_ldexp_f32_0_nopsub(float nofpclass(psub) %arg0) #0 {
define float @ret_ldexp_f32_0_nonsub(float nofpclass(nsub) %arg0) #0 {
; CHECK-LABEL: define nofpclass(nsub) float @ret_ldexp_f32_0_nonsub
; CHECK-SAME: (float nofpclass(nsub) [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nsub) float @llvm.ldexp.f32.i32(float nofpclass(nsub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nsub) float @llvm.ldexp.f32.i32(float nofpclass(nsub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -52,7 +53,7 @@ define float @ret_ldexp_f32_0_nonsub(float nofpclass(nsub) %arg0) #0 {
define float @ret_ldexp_f32_0_nosub(float nofpclass(sub) %arg0) #0 {
; CHECK-LABEL: define nofpclass(sub) float @ret_ldexp_f32_0_nosub
; CHECK-SAME: (float nofpclass(sub) [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) float @llvm.ldexp.f32.i32(float nofpclass(sub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) float @llvm.ldexp.f32.i32(float nofpclass(sub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -62,7 +63,7 @@ define float @ret_ldexp_f32_0_nosub(float nofpclass(sub) %arg0) #0 {
define float @ret_ldexp_f32_0_nosub_nosnan(float nofpclass(sub snan) %arg0) #0 {
; CHECK-LABEL: define nofpclass(snan sub) float @ret_ldexp_f32_0_nosub_nosnan
; CHECK-SAME: (float nofpclass(snan sub) [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(snan sub) float @llvm.ldexp.f32.i32(float nofpclass(snan sub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(snan sub) float @llvm.ldexp.f32.i32(float nofpclass(snan sub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -72,7 +73,7 @@ define float @ret_ldexp_f32_0_nosub_nosnan(float nofpclass(sub snan) %arg0) #0 {
define float @ret_ldexp_f32_0_nopsub_nopzero(float nofpclass(psub pzero) %arg0) #0 {
; CHECK-LABEL: define nofpclass(pzero psub) float @ret_ldexp_f32_0_nopsub_nopzero
; CHECK-SAME: (float nofpclass(pzero psub) [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pzero psub) float @llvm.ldexp.f32.i32(float nofpclass(pzero psub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pzero psub) float @llvm.ldexp.f32.i32(float nofpclass(pzero psub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -82,7 +83,7 @@ define float @ret_ldexp_f32_0_nopsub_nopzero(float nofpclass(psub pzero) %arg0)
define float @ret_ldexp_f32_0_nonsub_nonzero(float nofpclass(nsub nzero) %arg0) #0 {
; CHECK-LABEL: define nofpclass(nzero nsub) float @ret_ldexp_f32_0_nonsub_nonzero
; CHECK-SAME: (float nofpclass(nzero nsub) [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nzero nsub) float @llvm.ldexp.f32.i32(float nofpclass(nzero nsub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nzero nsub) float @llvm.ldexp.f32.i32(float nofpclass(nzero nsub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -92,7 +93,7 @@ define float @ret_ldexp_f32_0_nonsub_nonzero(float nofpclass(nsub nzero) %arg0)
define float @ret_ldexp_f32_0_nosub_nozero(float nofpclass(sub zero) %arg0) #0 {
; CHECK-LABEL: define nofpclass(zero sub) float @ret_ldexp_f32_0_nosub_nozero
; CHECK-SAME: (float nofpclass(zero sub) [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(zero sub) float @llvm.ldexp.f32.i32(float nofpclass(zero sub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(zero sub) float @llvm.ldexp.f32.i32(float nofpclass(zero sub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -102,7 +103,7 @@ define float @ret_ldexp_f32_0_nosub_nozero(float nofpclass(sub zero) %arg0) #0 {
define float @ret_ldexp_f32_0_nopsub_daz(float nofpclass(psub) %arg0) #1 {
; CHECK-LABEL: define nofpclass(psub) float @ret_ldexp_f32_0_nopsub_daz
; CHECK-SAME: (float nofpclass(psub) [[ARG0:%.*]]) #[[ATTR2:[0-9]+]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(psub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(psub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -112,7 +113,7 @@ define float @ret_ldexp_f32_0_nopsub_daz(float nofpclass(psub) %arg0) #1 {
define float @ret_ldexp_f32_0_nonsub_daz(float nofpclass(nsub) %arg0) #1 {
; CHECK-LABEL: define nofpclass(nsub) float @ret_ldexp_f32_0_nonsub_daz
; CHECK-SAME: (float nofpclass(nsub) [[ARG0:%.*]]) #[[ATTR2]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nsub) float @llvm.ldexp.f32.i32(float nofpclass(nsub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nsub) float @llvm.ldexp.f32.i32(float nofpclass(nsub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -122,7 +123,7 @@ define float @ret_ldexp_f32_0_nonsub_daz(float nofpclass(nsub) %arg0) #1 {
define float @ret_ldexp_f32_0_nosub_daz(float nofpclass(sub) %arg0) #1 {
; CHECK-LABEL: define nofpclass(sub) float @ret_ldexp_f32_0_nosub_daz
; CHECK-SAME: (float nofpclass(sub) [[ARG0:%.*]]) #[[ATTR2]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) float @llvm.ldexp.f32.i32(float nofpclass(sub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) float @llvm.ldexp.f32.i32(float nofpclass(sub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -132,7 +133,7 @@ define float @ret_ldexp_f32_0_nosub_daz(float nofpclass(sub) %arg0) #1 {
define float @ret_ldexp_f32_0_nosub_nosnan_daz(float nofpclass(sub snan) %arg0) #1 {
; CHECK-LABEL: define nofpclass(snan sub) float @ret_ldexp_f32_0_nosub_nosnan_daz
; CHECK-SAME: (float nofpclass(snan sub) [[ARG0:%.*]]) #[[ATTR2]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(snan sub) float @llvm.ldexp.f32.i32(float nofpclass(snan sub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(snan sub) float @llvm.ldexp.f32.i32(float nofpclass(snan sub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -142,7 +143,7 @@ define float @ret_ldexp_f32_0_nosub_nosnan_daz(float nofpclass(sub snan) %arg0)
define float @ret_ldexp_f32_0_nopsub_nozero_daz(float nofpclass(psub pzero) %arg0) #1 {
; CHECK-LABEL: define nofpclass(pzero psub) float @ret_ldexp_f32_0_nopsub_nozero_daz
; CHECK-SAME: (float nofpclass(pzero psub) [[ARG0:%.*]]) #[[ATTR2]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pzero psub) float @llvm.ldexp.f32.i32(float nofpclass(pzero psub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pzero psub) float @llvm.ldexp.f32.i32(float nofpclass(pzero psub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -152,7 +153,7 @@ define float @ret_ldexp_f32_0_nopsub_nozero_daz(float nofpclass(psub pzero) %arg
define float @ret_ldexp_f32_0_nonsub_nonzero_daz(float nofpclass(nsub nzero) %arg0) #1 {
; CHECK-LABEL: define nofpclass(nzero nsub) float @ret_ldexp_f32_0_nonsub_nonzero_daz
; CHECK-SAME: (float nofpclass(nzero nsub) [[ARG0:%.*]]) #[[ATTR2]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nzero nsub) float @llvm.ldexp.f32.i32(float nofpclass(nzero nsub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nzero nsub) float @llvm.ldexp.f32.i32(float nofpclass(nzero nsub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -162,7 +163,7 @@ define float @ret_ldexp_f32_0_nonsub_nonzero_daz(float nofpclass(nsub nzero) %ar
define float @ret_ldexp_f32_0_nosub_nozero_daz(float nofpclass(sub zero) %arg0) #1 {
; CHECK-LABEL: define nofpclass(zero sub) float @ret_ldexp_f32_0_nosub_nozero_daz
; CHECK-SAME: (float nofpclass(zero sub) [[ARG0:%.*]]) #[[ATTR2]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(zero sub) float @llvm.ldexp.f32.i32(float nofpclass(zero sub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(zero sub) float @llvm.ldexp.f32.i32(float nofpclass(zero sub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -172,7 +173,7 @@ define float @ret_ldexp_f32_0_nosub_nozero_daz(float nofpclass(sub zero) %arg0)
define float @ret_ldexp_f32_0_nopsub_dapz(float nofpclass(psub) %arg0) #2 {
; CHECK-LABEL: define nofpclass(psub) float @ret_ldexp_f32_0_nopsub_dapz
; CHECK-SAME: (float nofpclass(psub) [[ARG0:%.*]]) #[[ATTR3:[0-9]+]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(psub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(psub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -182,7 +183,7 @@ define float @ret_ldexp_f32_0_nopsub_dapz(float nofpclass(psub) %arg0) #2 {
define float @ret_ldexp_f32_0_nonsub_dapz(float nofpclass(nsub) %arg0) #2 {
; CHECK-LABEL: define nofpclass(nsub) float @ret_ldexp_f32_0_nonsub_dapz
; CHECK-SAME: (float nofpclass(nsub) [[ARG0:%.*]]) #[[ATTR3]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nsub) float @llvm.ldexp.f32.i32(float nofpclass(nsub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nsub) float @llvm.ldexp.f32.i32(float nofpclass(nsub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -192,7 +193,7 @@ define float @ret_ldexp_f32_0_nonsub_dapz(float nofpclass(nsub) %arg0) #2 {
define float @ret_ldexp_f32_0_nosub_dapz(float nofpclass(sub) %arg0) #2 {
; CHECK-LABEL: define nofpclass(sub) float @ret_ldexp_f32_0_nosub_dapz
; CHECK-SAME: (float nofpclass(sub) [[ARG0:%.*]]) #[[ATTR3]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) float @llvm.ldexp.f32.i32(float nofpclass(sub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) float @llvm.ldexp.f32.i32(float nofpclass(sub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -202,7 +203,7 @@ define float @ret_ldexp_f32_0_nosub_dapz(float nofpclass(sub) %arg0) #2 {
define float @ret_ldexp_f32_0_nopsub_nozero_dapz(float nofpclass(psub zero) %arg0) #2 {
; CHECK-LABEL: define nofpclass(nzero psub) float @ret_ldexp_f32_0_nopsub_nozero_dapz
; CHECK-SAME: (float nofpclass(zero psub) [[ARG0:%.*]]) #[[ATTR3]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nzero psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nzero psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -212,7 +213,7 @@ define float @ret_ldexp_f32_0_nopsub_nozero_dapz(float nofpclass(psub zero) %arg
define float @ret_ldexp_f32_0_nonsub_nonzero_dapz(float nofpclass(nsub nzero) %arg0) #2 {
; CHECK-LABEL: define nofpclass(nzero nsub) float @ret_ldexp_f32_0_nonsub_nonzero_dapz
; CHECK-SAME: (float nofpclass(nzero nsub) [[ARG0:%.*]]) #[[ATTR3]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nzero nsub) float @llvm.ldexp.f32.i32(float nofpclass(nzero nsub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nzero nsub) float @llvm.ldexp.f32.i32(float nofpclass(nzero nsub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -222,7 +223,7 @@ define float @ret_ldexp_f32_0_nonsub_nonzero_dapz(float nofpclass(nsub nzero) %a
define float @ret_ldexp_f32_0_nosub_nozero_dapz(float nofpclass(sub zero) %arg0) #2 {
; CHECK-LABEL: define nofpclass(zero sub) float @ret_ldexp_f32_0_nosub_nozero_dapz
; CHECK-SAME: (float nofpclass(zero sub) [[ARG0:%.*]]) #[[ATTR3]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(zero sub) float @llvm.ldexp.f32.i32(float nofpclass(zero sub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(zero sub) float @llvm.ldexp.f32.i32(float nofpclass(zero sub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -232,7 +233,7 @@ define float @ret_ldexp_f32_0_nosub_nozero_dapz(float nofpclass(sub zero) %arg0)
define float @ret_ldexp_f32_0_nopsub_nozero_dapz_dapz(float nofpclass(psub zero) %arg0) #2 {
; CHECK-LABEL: define nofpclass(nzero psub) float @ret_ldexp_f32_0_nopsub_nozero_dapz_dapz
; CHECK-SAME: (float nofpclass(zero psub) [[ARG0:%.*]]) #[[ATTR3]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nzero psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nzero psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -242,7 +243,7 @@ define float @ret_ldexp_f32_0_nopsub_nozero_dapz_dapz(float nofpclass(psub zero)
define float @ret_ldexp_f32_0_nopsub_nozero_daz_ieee(float nofpclass(psub zero) %arg0) #3 {
; CHECK-LABEL: define nofpclass(pzero psub) float @ret_ldexp_f32_0_nopsub_nozero_daz_ieee
; CHECK-SAME: (float nofpclass(zero psub) [[ARG0:%.*]]) #[[ATTR4:[0-9]+]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pzero psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pzero psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -252,7 +253,7 @@ define float @ret_ldexp_f32_0_nopsub_nozero_daz_ieee(float nofpclass(psub zero)
define float @ret_ldexp_f32_0_nopsub_nozero_daz_dynamic(float nofpclass(psub zero) %arg0) #4 {
; CHECK-LABEL: define nofpclass(psub) float @ret_ldexp_f32_0_nopsub_nozero_daz_dynamic
; CHECK-SAME: (float nofpclass(zero psub) [[ARG0:%.*]]) #[[ATTR5:[0-9]+]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -262,7 +263,7 @@ define float @ret_ldexp_f32_0_nopsub_nozero_daz_dynamic(float nofpclass(psub zer
define float @ret_ldexp_f32_0_nopsub_nozero_ieee_daz(float nofpclass(psub zero) %arg0) #5 {
; CHECK-LABEL: define nofpclass(pzero psub) float @ret_ldexp_f32_0_nopsub_nozero_ieee_daz
; CHECK-SAME: (float nofpclass(zero psub) [[ARG0:%.*]]) #[[ATTR6:[0-9]+]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pzero psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pzero psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -272,7 +273,7 @@ define float @ret_ldexp_f32_0_nopsub_nozero_ieee_daz(float nofpclass(psub zero)
define float @ret_ldexp_f32_0_nopsub_nozero_dynamic_daz(float nofpclass(psub zero) %arg0) #6 {
; CHECK-LABEL: define nofpclass(psub) float @ret_ldexp_f32_0_nopsub_nozero_dynamic_daz
; CHECK-SAME: (float nofpclass(zero psub) [[ARG0:%.*]]) #[[ATTR7:[0-9]+]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -282,7 +283,7 @@ define float @ret_ldexp_f32_0_nopsub_nozero_dynamic_daz(float nofpclass(psub zer
define float @ret_ldexp_f32_0_nopsub_nozero_dynamic_dapz(float nofpclass(psub zero) %arg0) #7 {
; CHECK-LABEL: define nofpclass(psub) float @ret_ldexp_f32_0_nopsub_nozero_dynamic_dapz
; CHECK-SAME: (float nofpclass(zero psub) [[ARG0:%.*]]) #[[ATTR8:[0-9]+]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -292,7 +293,7 @@ define float @ret_ldexp_f32_0_nopsub_nozero_dynamic_dapz(float nofpclass(psub ze
define float @ret_ldexp_f32_0_nopsub_nozero_dapz_dynamic(float nofpclass(psub zero) %arg0) #8 {
; CHECK-LABEL: define nofpclass(psub) float @ret_ldexp_f32_0_nopsub_nozero_dapz_dynamic
; CHECK-SAME: (float nofpclass(zero psub) [[ARG0:%.*]]) #[[ATTR9:[0-9]+]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -302,7 +303,7 @@ define float @ret_ldexp_f32_0_nopsub_nozero_dapz_dynamic(float nofpclass(psub ze
define float @ret_ldexp_f32_0_nopnorm(float nofpclass(pnorm) %arg0) #0 {
; CHECK-LABEL: define nofpclass(pnorm) float @ret_ldexp_f32_0_nopnorm
; CHECK-SAME: (float nofpclass(pnorm) [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pnorm) float @llvm.ldexp.f32.i32(float nofpclass(pnorm) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pnorm) float @llvm.ldexp.f32.i32(float nofpclass(pnorm) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -312,7 +313,7 @@ define float @ret_ldexp_f32_0_nopnorm(float nofpclass(pnorm) %arg0) #0 {
define float @ret_ldexp_f32_0_nnorm(float nofpclass(nnorm) %arg0) #0 {
; CHECK-LABEL: define nofpclass(nnorm) float @ret_ldexp_f32_0_nnorm
; CHECK-SAME: (float nofpclass(nnorm) [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nnorm) float @llvm.ldexp.f32.i32(float nofpclass(nnorm) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nnorm) float @llvm.ldexp.f32.i32(float nofpclass(nnorm) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -322,7 +323,7 @@ define float @ret_ldexp_f32_0_nnorm(float nofpclass(nnorm) %arg0) #0 {
define float @ret_ldexp_f32_0_nopnorm_nonsub(float nofpclass(pnorm nsub) %arg0) #0 {
; CHECK-LABEL: define nofpclass(nsub pnorm) float @ret_ldexp_f32_0_nopnorm_nonsub
; CHECK-SAME: (float nofpclass(nsub pnorm) [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nsub pnorm) float @llvm.ldexp.f32.i32(float nofpclass(nsub pnorm) [[ARG0]], i32 noundef 0) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nsub pnorm) float @llvm.ldexp.f32.i32(float nofpclass(nsub pnorm) [[ARG0]], i32 noundef 0) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 0)
@@ -332,7 +333,7 @@ define float @ret_ldexp_f32_0_nopnorm_nonsub(float nofpclass(pnorm nsub) %arg0)
define <2 x float> @ret_ldexp_v2f32_0(<2 x float> %arg0) #0 {
; CHECK-LABEL: define <2 x float> @ret_ldexp_v2f32_0
; CHECK-SAME: (<2 x float> [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call <2 x float> @llvm.ldexp.v2f32.v2i32(<2 x float> [[ARG0]], <2 x i32> <i32 undef, i32 0>) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call <2 x float> @llvm.ldexp.v2f32.v2i32(<2 x float> [[ARG0]], <2 x i32> <i32 undef, i32 0>) #[[ATTR12]]
; CHECK-NEXT: ret <2 x float> [[CALL]]
;
%call = call <2 x float> @llvm.ldexp.v2f32.v2i32(<2 x float> %arg0, <2 x i32> <i32 undef, i32 0>)
@@ -342,7 +343,7 @@ define <2 x float> @ret_ldexp_v2f32_0(<2 x float> %arg0) #0 {
define float @ret_ldexp_f32_i64(float %arg0, i64 %arg1) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_i64
; CHECK-SAME: (float [[ARG0:%.*]], i64 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i64(float [[ARG0]], i64 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i64(float [[ARG0]], i64 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i64(float %arg0, i64 %arg1)
@@ -352,7 +353,7 @@ define float @ret_ldexp_f32_i64(float %arg0, i64 %arg1) #0 {
define <2 x float> @ret_ldexp_v2f32(<2 x float> %arg0, <2 x i32> %arg1) #0 {
; CHECK-LABEL: define <2 x float> @ret_ldexp_v2f32
; CHECK-SAME: (<2 x float> [[ARG0:%.*]], <2 x i32> [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call <2 x float> @llvm.ldexp.v2f32.v2i32(<2 x float> [[ARG0]], <2 x i32> [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call <2 x float> @llvm.ldexp.v2f32.v2i32(<2 x float> [[ARG0]], <2 x i32> [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret <2 x float> [[CALL]]
;
%call = call <2 x float> @llvm.ldexp.v2f32.v2i32(<2 x float> %arg0, <2 x i32> %arg1)
@@ -362,7 +363,7 @@ define <2 x float> @ret_ldexp_v2f32(<2 x float> %arg0, <2 x i32> %arg1) #0 {
define float @ret_ldexp_f32_nonan(float nofpclass(nan) %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define nofpclass(nan) float @ret_ldexp_f32_nonan
; CHECK-SAME: (float nofpclass(nan) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nan) float @llvm.ldexp.f32.i32(float nofpclass(nan) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nan) float @llvm.ldexp.f32.i32(float nofpclass(nan) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -372,7 +373,7 @@ define float @ret_ldexp_f32_nonan(float nofpclass(nan) %arg0, i32 %arg1) #0 {
define float @ret_ldexp_f32_nosnan(float nofpclass(snan) %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define nofpclass(snan) float @ret_ldexp_f32_nosnan
; CHECK-SAME: (float nofpclass(snan) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(snan) float @llvm.ldexp.f32.i32(float nofpclass(snan) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(snan) float @llvm.ldexp.f32.i32(float nofpclass(snan) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -382,7 +383,7 @@ define float @ret_ldexp_f32_nosnan(float nofpclass(snan) %arg0, i32 %arg1) #0 {
define float @ret_ldexp_f32_noqnan(float nofpclass(qnan) %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_noqnan
; CHECK-SAME: (float nofpclass(qnan) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(qnan) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(qnan) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -392,7 +393,7 @@ define float @ret_ldexp_f32_noqnan(float nofpclass(qnan) %arg0, i32 %arg1) #0 {
define float @ret_ldexp_f32_noneg(float nofpclass(ninf nsub nnorm) %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define nofpclass(ninf nsub nnorm) float @ret_ldexp_f32_noneg
; CHECK-SAME: (float nofpclass(ninf nsub nnorm) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -402,7 +403,7 @@ define float @ret_ldexp_f32_noneg(float nofpclass(ninf nsub nnorm) %arg0, i32 %a
define float @ret_ldexp_f32_noneg_nonzero(float nofpclass(ninf nsub nnorm nzero) %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) float @ret_ldexp_f32_noneg_nonzero
; CHECK-SAME: (float nofpclass(ninf nzero nsub nnorm) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nzero nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nzero nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nzero nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nzero nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -412,7 +413,7 @@ define float @ret_ldexp_f32_noneg_nonzero(float nofpclass(ninf nsub nnorm nzero)
define float @ret_ldexp_f32_noneg_nozero(float nofpclass(ninf nsub nnorm) %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define nofpclass(ninf nsub nnorm) float @ret_ldexp_f32_noneg_nozero
; CHECK-SAME: (float nofpclass(ninf nsub nnorm) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -422,7 +423,7 @@ define float @ret_ldexp_f32_noneg_nozero(float nofpclass(ninf nsub nnorm) %arg0,
define float @ret_ldexp_f32_nonzero(float nofpclass(nzero) %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_nonzero
; CHECK-SAME: (float nofpclass(nzero) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(nzero) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(nzero) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -432,7 +433,7 @@ define float @ret_ldexp_f32_nonzero(float nofpclass(nzero) %arg0, i32 %arg1) #0
define float @ret_ldexp_f32_nopzero(float nofpclass(pzero) %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_nopzero
; CHECK-SAME: (float nofpclass(pzero) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(pzero) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(pzero) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -442,7 +443,7 @@ define float @ret_ldexp_f32_nopzero(float nofpclass(pzero) %arg0, i32 %arg1) #0
define float @ret_ldexp_f32_noneg_ftz_daz(float nofpclass(ninf nsub nnorm) %arg0, i32 %arg1) #1 {
; CHECK-LABEL: define nofpclass(ninf nsub nnorm) float @ret_ldexp_f32_noneg_ftz_daz
; CHECK-SAME: (float nofpclass(ninf nsub nnorm) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR2]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -452,7 +453,7 @@ define float @ret_ldexp_f32_noneg_ftz_daz(float nofpclass(ninf nsub nnorm) %arg0
define float @ret_ldexp_f32_noneg_nonzero_ftz_daz(float nofpclass(ninf nsub nnorm nzero) %arg0, i32 %arg1) #1 {
; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) float @ret_ldexp_f32_noneg_nonzero_ftz_daz
; CHECK-SAME: (float nofpclass(ninf nzero nsub nnorm) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR2]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nzero nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nzero nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nzero nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nzero nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -462,7 +463,7 @@ define float @ret_ldexp_f32_noneg_nonzero_ftz_daz(float nofpclass(ninf nsub nnor
define float @ret_ldexp_f32_noneg_nonzero_ftpz_dapz(float nofpclass(ninf nsub nnorm nzero) %arg0, i32 %arg1) #2 {
; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) float @ret_ldexp_f32_noneg_nonzero_ftpz_dapz
; CHECK-SAME: (float nofpclass(ninf nzero nsub nnorm) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR3]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nzero nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nzero nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nzero nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nzero nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -472,7 +473,7 @@ define float @ret_ldexp_f32_noneg_nonzero_ftpz_dapz(float nofpclass(ninf nsub nn
define float @ret_ldexp_f32_noninf_nonnorm(float nofpclass(ninf nnorm) %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_noninf_nonnorm
; CHECK-SAME: (float nofpclass(ninf nnorm) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(ninf nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(ninf nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -482,7 +483,7 @@ define float @ret_ldexp_f32_noninf_nonnorm(float nofpclass(ninf nnorm) %arg0, i3
define float @ret_ldexp_f32_noninf_nonnorm_ftz_daz(float nofpclass(ninf nnorm) %arg0, i32 %arg1) #1 {
; CHECK-LABEL: define float @ret_ldexp_f32_noninf_nonnorm_ftz_daz
; CHECK-SAME: (float nofpclass(ninf nnorm) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR2]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(ninf nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(ninf nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -492,7 +493,7 @@ define float @ret_ldexp_f32_noninf_nonnorm_ftz_daz(float nofpclass(ninf nnorm) %
define float @ret_ldexp_f32_noneg_ftz_ieee(float nofpclass(ninf nsub nnorm) %arg0, i32 %arg1) #3 {
; CHECK-LABEL: define nofpclass(ninf nsub nnorm) float @ret_ldexp_f32_noneg_ftz_ieee
; CHECK-SAME: (float nofpclass(ninf nsub nnorm) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR4]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -502,7 +503,7 @@ define float @ret_ldexp_f32_noneg_ftz_ieee(float nofpclass(ninf nsub nnorm) %arg
define float @ret_ldexp_f32_noneg_nonzero_ftz_ieee(float nofpclass(ninf nsub nnorm nzero) %arg0, i32 %arg1) #3 {
; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) float @ret_ldexp_f32_noneg_nonzero_ftz_ieee
; CHECK-SAME: (float nofpclass(ninf nzero nsub nnorm) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR4]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nzero nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nzero nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf nzero nsub nnorm) float @llvm.ldexp.f32.i32(float nofpclass(ninf nzero nsub nnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -512,7 +513,7 @@ define float @ret_ldexp_f32_noneg_nonzero_ftz_ieee(float nofpclass(ninf nsub nno
define float @ret_ldexp_f32_nopos(float nofpclass(pinf psub pnorm) %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define nofpclass(pinf psub pnorm) float @ret_ldexp_f32_nopos
; CHECK-SAME: (float nofpclass(pinf psub pnorm) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pinf psub pnorm) float @llvm.ldexp.f32.i32(float nofpclass(pinf psub pnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pinf psub pnorm) float @llvm.ldexp.f32.i32(float nofpclass(pinf psub pnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -522,7 +523,7 @@ define float @ret_ldexp_f32_nopos(float nofpclass(pinf psub pnorm) %arg0, i32 %a
define float @ret_ldexp_f32_nopos_nopzero(float nofpclass(pinf psub pnorm pzero) %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) float @ret_ldexp_f32_nopos_nopzero
; CHECK-SAME: (float nofpclass(pinf pzero psub pnorm) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pinf pzero psub pnorm) float @llvm.ldexp.f32.i32(float nofpclass(pinf pzero psub pnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pinf pzero psub pnorm) float @llvm.ldexp.f32.i32(float nofpclass(pinf pzero psub pnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -532,7 +533,7 @@ define float @ret_ldexp_f32_nopos_nopzero(float nofpclass(pinf psub pnorm pzero)
define float @ret_ldexp_f32_nopos_nozero(float nofpclass(pinf psub pnorm zero) %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) float @ret_ldexp_f32_nopos_nozero
; CHECK-SAME: (float nofpclass(pinf zero psub pnorm) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pinf pzero psub pnorm) float @llvm.ldexp.f32.i32(float nofpclass(pinf zero psub pnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pinf pzero psub pnorm) float @llvm.ldexp.f32.i32(float nofpclass(pinf zero psub pnorm) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -542,7 +543,7 @@ define float @ret_ldexp_f32_nopos_nozero(float nofpclass(pinf psub pnorm zero) %
define float @ret_ldexp_f32_nozero(float nofpclass(zero) %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_nozero
; CHECK-SAME: (float nofpclass(zero) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(zero) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(zero) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -552,7 +553,7 @@ define float @ret_ldexp_f32_nozero(float nofpclass(zero) %arg0, i32 %arg1) #0 {
define float @ret_ldexp_f32_noinf(float nofpclass(inf) %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_noinf
; CHECK-SAME: (float nofpclass(inf) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(inf) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(inf) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -562,7 +563,7 @@ define float @ret_ldexp_f32_noinf(float nofpclass(inf) %arg0, i32 %arg1) #0 {
define float @ret_ldexp_f32_noinf_nonan(float nofpclass(inf nan) %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define nofpclass(nan) float @ret_ldexp_f32_noinf_nonan
; CHECK-SAME: (float nofpclass(nan inf) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nan) float @llvm.ldexp.f32.i32(float nofpclass(nan inf) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nan) float @llvm.ldexp.f32.i32(float nofpclass(nan inf) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -573,7 +574,7 @@ define float @ret_ldexp_f32_known_pos_exp(float %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_known_pos_exp
; CHECK-SAME: (float [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[AND_ARG1:%.*]] = and i32 [[ARG1]], 127
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%and.arg1 = and i32 %arg1, 127
@@ -585,7 +586,7 @@ define float @ret_ldexp_f32_known_neg_exp(float %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_known_neg_exp
; CHECK-SAME: (float [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[OR_ARG1:%.*]] = or i32 [[ARG1]], -16
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%or.arg1 = or i32 %arg1, -16
@@ -597,7 +598,7 @@ define float @ret_ldexp_f32_known_pos_exp_nosub(float nofpclass(sub) %arg0, i32
; CHECK-LABEL: define nofpclass(sub) float @ret_ldexp_f32_known_pos_exp_nosub
; CHECK-SAME: (float nofpclass(sub) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[AND_ARG1:%.*]] = and i32 [[ARG1]], 127
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) float @llvm.ldexp.f32.i32(float nofpclass(sub) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) float @llvm.ldexp.f32.i32(float nofpclass(sub) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%and.arg1 = and i32 %arg1, 127
@@ -609,7 +610,7 @@ define float @ret_ldexp_f32_known_neg_exp_nosub(float nofpclass(sub) %arg0, i32
; CHECK-LABEL: define float @ret_ldexp_f32_known_neg_exp_nosub
; CHECK-SAME: (float nofpclass(sub) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[OR_ARG1:%.*]] = or i32 [[ARG1]], -16
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(sub) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(sub) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%or.arg1 = or i32 %arg1, -16
@@ -621,7 +622,7 @@ define float @ret_ldexp_f32_known_pos_exp_nopsub(float nofpclass(psub) %arg0, i3
; CHECK-LABEL: define nofpclass(psub) float @ret_ldexp_f32_known_pos_exp_nopsub
; CHECK-SAME: (float nofpclass(psub) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[AND_ARG1:%.*]] = and i32 [[ARG1]], 127
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(psub) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(psub) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%and.arg1 = and i32 %arg1, 127
@@ -633,7 +634,7 @@ define float @ret_ldexp_f32_known_neg_exp_nopsub(float nofpclass(psub) %arg0, i3
; CHECK-LABEL: define float @ret_ldexp_f32_known_neg_exp_nopsub
; CHECK-SAME: (float nofpclass(psub) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[OR_ARG1:%.*]] = or i32 [[ARG1]], -16
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(psub) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(psub) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%or.arg1 = or i32 %arg1, -16
@@ -645,7 +646,7 @@ define float @ret_ldexp_f32_known_pos_exp_nonsub(float nofpclass(nsub) %arg0, i3
; CHECK-LABEL: define nofpclass(nsub) float @ret_ldexp_f32_known_pos_exp_nonsub
; CHECK-SAME: (float nofpclass(nsub) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[AND_ARG1:%.*]] = and i32 [[ARG1]], 127
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nsub) float @llvm.ldexp.f32.i32(float nofpclass(nsub) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nsub) float @llvm.ldexp.f32.i32(float nofpclass(nsub) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%and.arg1 = and i32 %arg1, 127
@@ -657,7 +658,7 @@ define float @ret_ldexp_f32_known_neg_exp_nonsub(float nofpclass(nsub) %arg0, i3
; CHECK-LABEL: define float @ret_ldexp_f32_known_neg_exp_nonsub
; CHECK-SAME: (float nofpclass(nsub) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[OR_ARG1:%.*]] = or i32 [[ARG1]], -16
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(nsub) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(nsub) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%or.arg1 = or i32 %arg1, -16
@@ -669,7 +670,7 @@ define float @ret_ldexp_f32_known_pos_exp_noinf(float nofpclass(inf) %arg0, i32
; CHECK-LABEL: define float @ret_ldexp_f32_known_pos_exp_noinf
; CHECK-SAME: (float nofpclass(inf) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[AND_ARG1:%.*]] = and i32 [[ARG1]], 127
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(inf) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(inf) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%and.arg1 = and i32 %arg1, 127
@@ -681,7 +682,7 @@ define float @ret_ldexp_f32_known_neg_exp_noinf(float nofpclass(inf) %arg0, i32
; CHECK-LABEL: define nofpclass(inf) float @ret_ldexp_f32_known_neg_exp_noinf
; CHECK-SAME: (float nofpclass(inf) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[OR_ARG1:%.*]] = or i32 [[ARG1]], -16
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(inf) float @llvm.ldexp.f32.i32(float nofpclass(inf) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(inf) float @llvm.ldexp.f32.i32(float nofpclass(inf) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%or.arg1 = or i32 %arg1, -16
@@ -693,7 +694,7 @@ define <2 x float> @ret_ldexp_v2f32_known_pos_exp_noinf(<2 x float> nofpclass(in
; CHECK-LABEL: define <2 x float> @ret_ldexp_v2f32_known_pos_exp_noinf
; CHECK-SAME: (<2 x float> nofpclass(inf) [[ARG0:%.*]], <2 x i32> [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[AND_ARG1:%.*]] = and <2 x i32> [[ARG1]], <i32 127, i32 64>
-; CHECK-NEXT: [[CALL:%.*]] = call <2 x float> @llvm.ldexp.v2f32.v2i32(<2 x float> nofpclass(inf) [[ARG0]], <2 x i32> [[AND_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call <2 x float> @llvm.ldexp.v2f32.v2i32(<2 x float> nofpclass(inf) [[ARG0]], <2 x i32> [[AND_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret <2 x float> [[CALL]]
;
%and.arg1 = and <2 x i32> %arg1, <i32 127, i32 64>
@@ -705,7 +706,7 @@ define <2 x float> @ret_ldexp_v2f32_known_neg_exp_noinf_splat(<2 x float> nofpcl
; CHECK-LABEL: define nofpclass(inf) <2 x float> @ret_ldexp_v2f32_known_neg_exp_noinf_splat
; CHECK-SAME: (<2 x float> nofpclass(inf) [[ARG0:%.*]], <2 x i32> [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[OR_ARG1:%.*]] = or <2 x i32> [[ARG1]], splat (i32 -32)
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(inf) <2 x float> @llvm.ldexp.v2f32.v2i32(<2 x float> nofpclass(inf) [[ARG0]], <2 x i32> [[OR_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(inf) <2 x float> @llvm.ldexp.v2f32.v2i32(<2 x float> nofpclass(inf) [[ARG0]], <2 x i32> [[OR_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret <2 x float> [[CALL]]
;
%or.arg1 = or <2 x i32> %arg1, splat (i32 -32)
@@ -718,7 +719,7 @@ define <2 x float> @ret_ldexp_v2f32_known_neg_exp_noinf_nonsplat(<2 x float> nof
; CHECK-LABEL: define <2 x float> @ret_ldexp_v2f32_known_neg_exp_noinf_nonsplat
; CHECK-SAME: (<2 x float> nofpclass(inf) [[ARG0:%.*]], <2 x i32> [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[OR_ARG1:%.*]] = or <2 x i32> [[ARG1]], <i32 -16, i32 -32>
-; CHECK-NEXT: [[CALL:%.*]] = call <2 x float> @llvm.ldexp.v2f32.v2i32(<2 x float> nofpclass(inf) [[ARG0]], <2 x i32> [[OR_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call <2 x float> @llvm.ldexp.v2f32.v2i32(<2 x float> nofpclass(inf) [[ARG0]], <2 x i32> [[OR_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret <2 x float> [[CALL]]
;
%or.arg1 = or <2 x i32> %arg1, <i32 -16, i32 -32>
@@ -730,7 +731,7 @@ define float @ret_ldexp_f32_known_pos_exp_nopinf(float nofpclass(pinf) %arg0, i3
; CHECK-LABEL: define float @ret_ldexp_f32_known_pos_exp_nopinf
; CHECK-SAME: (float nofpclass(pinf) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[AND_ARG1:%.*]] = and i32 [[ARG1]], 127
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(pinf) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(pinf) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%and.arg1 = and i32 %arg1, 127
@@ -742,7 +743,7 @@ define float @ret_ldexp_f32_known_neg_exp_nopinf(float nofpclass(pinf) %arg0, i3
; CHECK-LABEL: define nofpclass(pinf) float @ret_ldexp_f32_known_neg_exp_nopinf
; CHECK-SAME: (float nofpclass(pinf) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[OR_ARG1:%.*]] = or i32 [[ARG1]], -16
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pinf) float @llvm.ldexp.f32.i32(float nofpclass(pinf) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pinf) float @llvm.ldexp.f32.i32(float nofpclass(pinf) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%or.arg1 = or i32 %arg1, -16
@@ -754,7 +755,7 @@ define float @ret_ldexp_f32_known_pos_exp_noninf(float nofpclass(ninf) %arg0, i3
; CHECK-LABEL: define float @ret_ldexp_f32_known_pos_exp_noninf
; CHECK-SAME: (float nofpclass(ninf) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[AND_ARG1:%.*]] = and i32 [[ARG1]], 127
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(ninf) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(ninf) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%and.arg1 = and i32 %arg1, 127
@@ -766,7 +767,7 @@ define float @ret_ldexp_f32_known_neg_exp_noninf(float nofpclass(ninf) %arg0, i3
; CHECK-LABEL: define nofpclass(ninf) float @ret_ldexp_f32_known_neg_exp_noninf
; CHECK-SAME: (float nofpclass(ninf) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
; CHECK-NEXT: [[OR_ARG1:%.*]] = or i32 [[ARG1]], -16
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf) float @llvm.ldexp.f32.i32(float nofpclass(ninf) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(ninf) float @llvm.ldexp.f32.i32(float nofpclass(ninf) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%or.arg1 = or i32 %arg1, -16
@@ -778,7 +779,7 @@ define float @ret_ldexp_f32_known_pos_exp_nopsub_nopzero_daz(float nofpclass(psu
; CHECK-LABEL: define nofpclass(pzero psub) float @ret_ldexp_f32_known_pos_exp_nopsub_nopzero_daz
; CHECK-SAME: (float nofpclass(pzero psub) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR2]] {
; CHECK-NEXT: [[AND_ARG1:%.*]] = and i32 [[ARG1]], 127
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pzero psub) float @llvm.ldexp.f32.i32(float nofpclass(pzero psub) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pzero psub) float @llvm.ldexp.f32.i32(float nofpclass(pzero psub) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%and.arg1 = and i32 %arg1, 127
@@ -790,7 +791,7 @@ define float @ret_ldexp_f32_known_neg_exp_nopsub_nopzero_daz(float nofpclass(psu
; CHECK-LABEL: define float @ret_ldexp_f32_known_neg_exp_nopsub_nopzero_daz
; CHECK-SAME: (float nofpclass(pzero psub) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR2]] {
; CHECK-NEXT: [[OR_ARG1:%.*]] = or i32 [[ARG1]], -16
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(pzero psub) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(pzero psub) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%or.arg1 = or i32 %arg1, -16
@@ -802,7 +803,7 @@ define float @ret_ldexp_f32_known_pos_exp_nonsub_nonzero_daz(float nofpclass(nsu
; CHECK-LABEL: define nofpclass(nzero nsub) float @ret_ldexp_f32_known_pos_exp_nonsub_nonzero_daz
; CHECK-SAME: (float nofpclass(nzero nsub) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR2]] {
; CHECK-NEXT: [[AND_ARG1:%.*]] = and i32 [[ARG1]], 127
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nzero nsub) float @llvm.ldexp.f32.i32(float nofpclass(nzero nsub) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nzero nsub) float @llvm.ldexp.f32.i32(float nofpclass(nzero nsub) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%and.arg1 = and i32 %arg1, 127
@@ -814,7 +815,7 @@ define float @ret_ldexp_f32_known_neg_exp_nonsub_nonzero_daz(float nofpclass(nsu
; CHECK-LABEL: define float @ret_ldexp_f32_known_neg_exp_nonsub_nonzero_daz
; CHECK-SAME: (float nofpclass(nzero nsub) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR2]] {
; CHECK-NEXT: [[OR_ARG1:%.*]] = or i32 [[ARG1]], -16
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(nzero nsub) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(nzero nsub) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%or.arg1 = or i32 %arg1, -16
@@ -826,7 +827,7 @@ define float @ret_ldexp_f32_known_pos_exp_nosub_nozero_daz(float nofpclass(sub z
; CHECK-LABEL: define nofpclass(zero sub) float @ret_ldexp_f32_known_pos_exp_nosub_nozero_daz
; CHECK-SAME: (float nofpclass(zero sub) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR2]] {
; CHECK-NEXT: [[AND_ARG1:%.*]] = and i32 [[ARG1]], 127
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(zero sub) float @llvm.ldexp.f32.i32(float nofpclass(zero sub) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(zero sub) float @llvm.ldexp.f32.i32(float nofpclass(zero sub) [[ARG0]], i32 [[AND_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%and.arg1 = and i32 %arg1, 127
@@ -838,7 +839,7 @@ define float @ret_ldexp_f32_known_neg_exp_nosub_nozero_daz(float nofpclass(sub z
; CHECK-LABEL: define float @ret_ldexp_f32_known_neg_exp_nosub_nozero_daz
; CHECK-SAME: (float nofpclass(zero sub) [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR2]] {
; CHECK-NEXT: [[OR_ARG1:%.*]] = or i32 [[ARG1]], -16
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(zero sub) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(zero sub) [[ARG0]], i32 [[OR_ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%or.arg1 = or i32 %arg1, -16
@@ -849,7 +850,7 @@ define float @ret_ldexp_f32_known_neg_exp_nosub_nozero_daz(float nofpclass(sub z
define float @ret_ldexp_f32_22(float %arg0) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_22
; CHECK-SAME: (float [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef 22) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef 22) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 22)
@@ -859,7 +860,7 @@ define float @ret_ldexp_f32_22(float %arg0) #0 {
define float @ret_ldexp_f32_23(float %arg0) #0 {
; CHECK-LABEL: define nofpclass(sub) float @ret_ldexp_f32_23
; CHECK-SAME: (float [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef 23) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef 23) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 23)
@@ -869,7 +870,7 @@ define float @ret_ldexp_f32_23(float %arg0) #0 {
define float @ret_ldexp_f32_24(float %arg0) #0 {
; CHECK-LABEL: define nofpclass(sub) float @ret_ldexp_f32_24
; CHECK-SAME: (float [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef 24) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef 24) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 24)
@@ -879,8 +880,8 @@ define float @ret_ldexp_f32_24(float %arg0) #0 {
define float @ret_ldexp_f32_min24(float %arg0, i32 %arg1) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_min24
; CHECK-SAME: (float [[ARG0:%.*]], i32 [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[MIN:%.*]] = call i32 @llvm.smin.i32(i32 [[ARG1]], i32 noundef 24) #[[ATTR10]]
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 [[MIN]]) #[[ATTR10]]
+; CHECK-NEXT: [[MIN:%.*]] = call i32 @llvm.smin.i32(i32 [[ARG1]], i32 noundef 24) #[[ATTR12]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 [[MIN]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%min = call i32 @llvm.smin.i32(i32 %arg1, i32 24)
@@ -891,7 +892,7 @@ define float @ret_ldexp_f32_min24(float %arg0, i32 %arg1) #0 {
define float @ret_ldexp_f32_23_nnan(float nofpclass(nan) %arg0) #0 {
; CHECK-LABEL: define nofpclass(nan sub) float @ret_ldexp_f32_23_nnan
; CHECK-SAME: (float nofpclass(nan) [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nan sub) float @llvm.ldexp.f32.i32(float nofpclass(nan) [[ARG0]], i32 noundef 23) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nan sub) float @llvm.ldexp.f32.i32(float nofpclass(nan) [[ARG0]], i32 noundef 23) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 23)
@@ -901,7 +902,7 @@ define float @ret_ldexp_f32_23_nnan(float nofpclass(nan) %arg0) #0 {
define double @ret_ldexp_f64_24(double %arg0) #0 {
; CHECK-LABEL: define double @ret_ldexp_f64_24
; CHECK-SAME: (double [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call double @llvm.ldexp.f64.i32(double [[ARG0]], i32 noundef 24) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call double @llvm.ldexp.f64.i32(double [[ARG0]], i32 noundef 24) #[[ATTR12]]
; CHECK-NEXT: ret double [[CALL]]
;
%call = call double @llvm.ldexp.f64.i32(double %arg0, i32 24)
@@ -911,7 +912,7 @@ define double @ret_ldexp_f64_24(double %arg0) #0 {
define double @ret_ldexp_f64_51(double %arg0) #0 {
; CHECK-LABEL: define double @ret_ldexp_f64_51
; CHECK-SAME: (double [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call double @llvm.ldexp.f64.i32(double [[ARG0]], i32 noundef 51) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call double @llvm.ldexp.f64.i32(double [[ARG0]], i32 noundef 51) #[[ATTR12]]
; CHECK-NEXT: ret double [[CALL]]
;
%call = call double @llvm.ldexp.f64.i32(double %arg0, i32 51)
@@ -921,7 +922,7 @@ define double @ret_ldexp_f64_51(double %arg0) #0 {
define double @ret_ldexp_f64_52(double %arg0) #0 {
; CHECK-LABEL: define nofpclass(sub) double @ret_ldexp_f64_52
; CHECK-SAME: (double [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) double @llvm.ldexp.f64.i32(double [[ARG0]], i32 noundef 52) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) double @llvm.ldexp.f64.i32(double [[ARG0]], i32 noundef 52) #[[ATTR12]]
; CHECK-NEXT: ret double [[CALL]]
;
%call = call double @llvm.ldexp.f64.i32(double %arg0, i32 52)
@@ -931,7 +932,7 @@ define double @ret_ldexp_f64_52(double %arg0) #0 {
define double @ret_ldexp_f64_53(double %arg0) #0 {
; CHECK-LABEL: define nofpclass(sub) double @ret_ldexp_f64_53
; CHECK-SAME: (double [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) double @llvm.ldexp.f64.i32(double [[ARG0]], i32 noundef 53) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) double @llvm.ldexp.f64.i32(double [[ARG0]], i32 noundef 53) #[[ATTR12]]
; CHECK-NEXT: ret double [[CALL]]
;
%call = call double @llvm.ldexp.f64.i32(double %arg0, i32 53)
@@ -941,7 +942,7 @@ define double @ret_ldexp_f64_53(double %arg0) #0 {
define half @ret_ldexp_f16_8(half %arg0) #0 {
; CHECK-LABEL: define half @ret_ldexp_f16_8
; CHECK-SAME: (half [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call half @llvm.ldexp.f16.i32(half [[ARG0]], i32 noundef 8) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call half @llvm.ldexp.f16.i32(half [[ARG0]], i32 noundef 8) #[[ATTR12]]
; CHECK-NEXT: ret half [[CALL]]
;
%call = call half @llvm.ldexp.f16.i32(half %arg0, i32 8)
@@ -951,7 +952,7 @@ define half @ret_ldexp_f16_8(half %arg0) #0 {
define half @ret_ldexp_f16_9(half %arg0) #0 {
; CHECK-LABEL: define half @ret_ldexp_f16_9
; CHECK-SAME: (half [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call half @llvm.ldexp.f16.i32(half [[ARG0]], i32 noundef 9) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call half @llvm.ldexp.f16.i32(half [[ARG0]], i32 noundef 9) #[[ATTR12]]
; CHECK-NEXT: ret half [[CALL]]
;
%call = call half @llvm.ldexp.f16.i32(half %arg0, i32 9)
@@ -961,7 +962,7 @@ define half @ret_ldexp_f16_9(half %arg0) #0 {
define half @ret_ldexp_f16_10(half %arg0) #0 {
; CHECK-LABEL: define nofpclass(sub) half @ret_ldexp_f16_10
; CHECK-SAME: (half [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) half @llvm.ldexp.f16.i32(half [[ARG0]], i32 noundef 10) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) half @llvm.ldexp.f16.i32(half [[ARG0]], i32 noundef 10) #[[ATTR12]]
; CHECK-NEXT: ret half [[CALL]]
;
%call = call half @llvm.ldexp.f16.i32(half %arg0, i32 10)
@@ -971,7 +972,7 @@ define half @ret_ldexp_f16_10(half %arg0) #0 {
define bfloat @ret_ldexp_bf16_6(bfloat %arg0) #0 {
; CHECK-LABEL: define bfloat @ret_ldexp_bf16_6
; CHECK-SAME: (bfloat [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call bfloat @llvm.ldexp.bf16.i32(bfloat [[ARG0]], i32 noundef 6) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call bfloat @llvm.ldexp.bf16.i32(bfloat [[ARG0]], i32 noundef 6) #[[ATTR12]]
; CHECK-NEXT: ret bfloat [[CALL]]
;
%call = call bfloat @llvm.ldexp.bf16.i32(bfloat %arg0, i32 6)
@@ -981,7 +982,7 @@ define bfloat @ret_ldexp_bf16_6(bfloat %arg0) #0 {
define bfloat @ret_ldexp_bf16_7(bfloat %arg0) #0 {
; CHECK-LABEL: define nofpclass(sub) bfloat @ret_ldexp_bf16_7
; CHECK-SAME: (bfloat [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) bfloat @llvm.ldexp.bf16.i32(bfloat [[ARG0]], i32 noundef 7) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) bfloat @llvm.ldexp.bf16.i32(bfloat [[ARG0]], i32 noundef 7) #[[ATTR12]]
; CHECK-NEXT: ret bfloat [[CALL]]
;
%call = call bfloat @llvm.ldexp.bf16.i32(bfloat %arg0, i32 7)
@@ -991,7 +992,7 @@ define bfloat @ret_ldexp_bf16_7(bfloat %arg0) #0 {
define bfloat @ret_ldexp_bf16_8(bfloat %arg0) #0 {
; CHECK-LABEL: define nofpclass(sub) bfloat @ret_ldexp_bf16_8
; CHECK-SAME: (bfloat [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) bfloat @llvm.ldexp.bf16.i32(bfloat [[ARG0]], i32 noundef 8) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(sub) bfloat @llvm.ldexp.bf16.i32(bfloat [[ARG0]], i32 noundef 8) #[[ATTR12]]
; CHECK-NEXT: ret bfloat [[CALL]]
;
%call = call bfloat @llvm.ldexp.bf16.i32(bfloat %arg0, i32 8)
@@ -1001,7 +1002,7 @@ define bfloat @ret_ldexp_bf16_8(bfloat %arg0) #0 {
define float @ret_ldexp_f32_neg22(float %arg0) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_neg22
; CHECK-SAME: (float [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef -22) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef -22) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 -22)
@@ -1011,7 +1012,7 @@ define float @ret_ldexp_f32_neg22(float %arg0) #0 {
define float @ret_ldexp_f32_neg23(float %arg0) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_neg23
; CHECK-SAME: (float [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef -23) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef -23) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 -23)
@@ -1021,7 +1022,7 @@ define float @ret_ldexp_f32_neg23(float %arg0) #0 {
define float @ret_ldexp_f32_neg24(float %arg0) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_neg24
; CHECK-SAME: (float [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef -24) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef -24) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 -24)
@@ -1031,7 +1032,7 @@ define float @ret_ldexp_f32_neg24(float %arg0) #0 {
define float @ret_ldexp_f32_neg126(float %arg0) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_neg126
; CHECK-SAME: (float [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef -126) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef -126) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 -126)
@@ -1041,7 +1042,7 @@ define float @ret_ldexp_f32_neg126(float %arg0) #0 {
define float @ret_ldexp_f32_neg127(float %arg0) #0 {
; CHECK-LABEL: define float @ret_ldexp_f32_neg127
; CHECK-SAME: (float [[ARG0:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef -127) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float [[ARG0]], i32 noundef -127) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 -127)
@@ -1051,7 +1052,119 @@ define float @ret_ldexp_f32_neg127(float %arg0) #0 {
define float @ret_ldexp_f32_noinf_exp_known_neg_or_0(float nofpclass(inf) %arg0, i32 range(i32 -256, 1) %arg1) #0 {
; CHECK-LABEL: define nofpclass(inf) float @ret_ldexp_f32_noinf_exp_known_neg_or_0
; CHECK-SAME: (float nofpclass(inf) [[ARG0:%.*]], i32 range(i32 -256, 1) [[ARG1:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(inf) float @llvm.ldexp.f32.i32(float nofpclass(inf) [[ARG0]], i32 [[ARG1]]) #[[ATTR10]]
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(inf) float @llvm.ldexp.f32.i32(float nofpclass(inf) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
+; CHECK-NEXT: ret float [[CALL]]
+;
+ %call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
+ ret float %call
+}
+
+; Negative test, a subnormal result may still be flushed to zero on output.
+define float @ret_ldexp_f32_nozero_exp_known_nonneg_ftz_ieee(float nofpclass(zero) %arg0, i32 range(i32 0, 256) %arg1) #3 {
+; CHECK-LABEL: define float @ret_ldexp_f32_nozero_exp_known_nonneg_ftz_ieee
+; CHECK-SAME: (float nofpclass(zero) [[ARG0:%.*]], i32 range(i32 0, 256) [[ARG1:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[CALL:%.*]] = call float @llvm.ldexp.f32.i32(float nofpclass(zero) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
+; CHECK-NEXT: ret float [[CALL]]
+;
+ %call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
+ ret float %call
+}
+
+; The exponent scales but never negates. Only the positive side is ruled out.
+define float @ret_ldexp_f32_nopsub_nopzero_exp_known_nonneg(float nofpclass(psub pzero) %arg0, i32 range(i32 0, 256) %arg1) #0 {
+; CHECK-LABEL: define nofpclass(pzero psub) float @ret_ldexp_f32_nopsub_nopzero_exp_known_nonneg
+; CHECK-SAME: (float nofpclass(pzero psub) [[ARG0:%.*]], i32 range(i32 0, 256) [[ARG1:%.*]]) #[[ATTR1]] {
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pzero psub) float @llvm.ldexp.f32.i32(float nofpclass(pzero psub) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
+; CHECK-NEXT: ret float [[CALL]]
+;
+ %call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
+ ret float %call
+}
+
+; A source already free of zero, subnormal and inf still learns from the
+; exponent. A non-negative exponent cannot underflow.
+define float @ret_ldexp_f32_nozerosubinf_exp_known_nonneg(float nofpclass(zero sub inf) %arg0, i32 range(i32 0, 256) %arg1) #0 {
+; CHECK-LABEL: define nofpclass(zero sub) float @ret_ldexp_f32_nozerosubinf_exp_known_nonneg
+; CHECK-SAME: (float nofpclass(inf zero sub) [[ARG0:%.*]], i32 range(i32 0, 256) [[ARG1:%.*]]) #[[ATTR1]] {
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(zero sub) float @llvm.ldexp.f32.i32(float nofpclass(inf zero sub) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
+; CHECK-NEXT: ret float [[CALL]]
+;
+ %call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
+ ret float %call
+}
+
+; An exponent wider than int saturates rather than truncating. A full i64
+; range stays unconstrained and the result may still overflow.
+define float @ret_ldexp_f32_noinfnan_exp_i64(float nofpclass(inf nan) %arg0, i64 %arg1) #0 {
+; CHECK-LABEL: define nofpclass(nan) float @ret_ldexp_f32_noinfnan_exp_i64
+; CHECK-SAME: (float nofpclass(nan inf) [[ARG0:%.*]], i64 [[ARG1:%.*]]) #[[ATTR1]] {
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nan) float @llvm.ldexp.f32.i64(float nofpclass(nan inf) [[ARG0]], i64 [[ARG1]]) #[[ATTR12]]
+; CHECK-NEXT: ret float [[CALL]]
+;
+ %call = call float @llvm.ldexp.f32.i64(float %arg0, i64 %arg1)
+ ret float %call
+}
+
+; A non-positive exponent cannot grow the result. A normal or an infinity
+; must come from a source of the same class.
+define float @ret_ldexp_f32_nonorminf_exp_nonpos(float nofpclass(norm inf) %arg0, i32 range(i32 -256, 1) %arg1) #0 {
+; CHECK-LABEL: define nofpclass(inf norm) float @ret_ldexp_f32_nonorminf_exp_nonpos
+; CHECK-SAME: (float nofpclass(inf norm) [[ARG0:%.*]], i32 range(i32 -256, 1) [[ARG1:%.*]]) #[[ATTR1]] {
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(inf norm) float @llvm.ldexp.f32.i32(float nofpclass(inf norm) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
+; CHECK-NEXT: ret float [[CALL]]
+;
+ %call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
+ ret float %call
+}
+
+; Negative test, double-double reports no precision, so no exponent is known to
+; leave the subnormal range.
+define ppc_fp128 @ret_ldexp_ppcf128_exp_known_nonneg(ppc_fp128 %arg0, i32 range(i32 0, 256) %arg1) #0 {
+; CHECK-LABEL: define ppc_fp128 @ret_ldexp_ppcf128_exp_known_nonneg
+; CHECK-SAME: (ppc_fp128 [[ARG0:%.*]], i32 range(i32 0, 256) [[ARG1:%.*]]) #[[ATTR1]] {
+; CHECK-NEXT: [[CALL:%.*]] = call ppc_fp128 @llvm.ldexp.ppcf128.i32(ppc_fp128 [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
+; CHECK-NEXT: ret ppc_fp128 [[CALL]]
+;
+ %call = call ppc_fp128 @llvm.ldexp.ppcf128.i32(ppc_fp128 %arg0, i32 %arg1)
+ ret ppc_fp128 %call
+}
+
+; The output mode decides which zero a flushed subnormal can reach. A source
+; that rules out one sign of subnormal separates all four modes.
+define float @ret_ldexp_f32_nozero_nopsub_exp_nonneg_ieee_out(float nofpclass(zero psub) %arg0, i32 range(i32 0, 256) %arg1) #0 {
+; CHECK-LABEL: define nofpclass(zero psub) float @ret_ldexp_f32_nozero_nopsub_exp_nonneg_ieee_out
+; CHECK-SAME: (float nofpclass(zero psub) [[ARG0:%.*]], i32 range(i32 0, 256) [[ARG1:%.*]]) #[[ATTR1]] {
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(zero psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
+; CHECK-NEXT: ret float [[CALL]]
+;
+ %call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
+ ret float %call
+}
+
+define float @ret_ldexp_f32_nozero_nopsub_exp_nonneg_preservesign_out(float nofpclass(zero psub) %arg0, i32 range(i32 0, 256) %arg1) #3 {
+; CHECK-LABEL: define nofpclass(pzero psub) float @ret_ldexp_f32_nozero_nopsub_exp_nonneg_preservesign_out
+; CHECK-SAME: (float nofpclass(zero psub) [[ARG0:%.*]], i32 range(i32 0, 256) [[ARG1:%.*]]) #[[ATTR4]] {
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(pzero psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
+; CHECK-NEXT: ret float [[CALL]]
+;
+ %call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
+ ret float %call
+}
+
+define float @ret_ldexp_f32_nozero_nopsub_exp_nonneg_positivezero_out(float nofpclass(zero psub) %arg0, i32 range(i32 0, 256) %arg1) #9 {
+; CHECK-LABEL: define nofpclass(nzero psub) float @ret_ldexp_f32_nozero_nopsub_exp_nonneg_positivezero_out
+; CHECK-SAME: (float nofpclass(zero psub) [[ARG0:%.*]], i32 range(i32 0, 256) [[ARG1:%.*]]) #[[ATTR10:[0-9]+]] {
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(nzero psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
+; CHECK-NEXT: ret float [[CALL]]
+;
+ %call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
+ ret float %call
+}
+
+define float @ret_ldexp_f32_nozero_nopsub_exp_nonneg_dynamic_out(float nofpclass(zero psub) %arg0, i32 range(i32 0, 256) %arg1) #10 {
+; CHECK-LABEL: define nofpclass(psub) float @ret_ldexp_f32_nozero_nopsub_exp_nonneg_dynamic_out
+; CHECK-SAME: (float nofpclass(zero psub) [[ARG0:%.*]], i32 range(i32 0, 256) [[ARG1:%.*]]) #[[ATTR11:[0-9]+]] {
+; CHECK-NEXT: [[CALL:%.*]] = call nofpclass(psub) float @llvm.ldexp.f32.i32(float nofpclass(zero psub) [[ARG0]], i32 [[ARG1]]) #[[ATTR12]]
; CHECK-NEXT: ret float [[CALL]]
;
%call = call float @llvm.ldexp.f32.i32(float %arg0, i32 %arg1)
@@ -1067,6 +1180,8 @@ attributes #5 = { denormal_fpenv(ieee|preservesign) }
attributes #6 = { denormal_fpenv(dynamic|preservesign) }
attributes #7 = { denormal_fpenv(dynamic|positivezero) }
attributes #8 = { denormal_fpenv(positivezero|dynamic) }
+attributes #9 = { denormal_fpenv(positivezero|ieee) }
+attributes #10 = { denormal_fpenv(dynamic|ieee) }
;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
; TUNIT: {{.*}}
diff --git a/llvm/test/Transforms/Attributor/nofpclass.ll b/llvm/test/Transforms/Attributor/nofpclass.ll
index 719896870a479..96a3ee4bda508 100644
--- a/llvm/test/Transforms/Attributor/nofpclass.ll
+++ b/llvm/test/Transforms/Attributor/nofpclass.ll
@@ -3163,7 +3163,7 @@ define float @fadd_double_no_noninf_nnorm(float noundef nofpclass(ninf nnorm) %a
define float @fadd_double_no_pnorm_psub(float noundef nofpclass(pnorm psub) %arg) {
; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
-; CHECK-LABEL: define noundef float @fadd_double_no_pnorm_psub
+; CHECK-LABEL: define noundef nofpclass(psub) float @fadd_double_no_pnorm_psub
; CHECK-SAME: (float noundef nofpclass(psub pnorm) [[ARG:%.*]]) #[[ATTR3]] {
; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
; CHECK-NEXT: ret float [[ADD]]
@@ -3174,7 +3174,7 @@ define float @fadd_double_no_pnorm_psub(float noundef nofpclass(pnorm psub) %arg
define float @fadd_double_no_nnorm_nsub(float noundef nofpclass(nnorm nsub) %arg) {
; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
-; CHECK-LABEL: define noundef float @fadd_double_no_nnorm_nsub
+; CHECK-LABEL: define noundef nofpclass(nsub) float @fadd_double_no_nnorm_nsub
; CHECK-SAME: (float noundef nofpclass(nsub nnorm) [[ARG:%.*]]) #[[ATTR3]] {
; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
; CHECK-NEXT: ret float [[ADD]]
@@ -3185,7 +3185,7 @@ define float @fadd_double_no_nnorm_nsub(float noundef nofpclass(nnorm nsub) %arg
define float @fadd_double_no_nopsub_pzero(float noundef nofpclass(psub pzero) %arg) {
; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
-; CHECK-LABEL: define noundef nofpclass(pzero) float @fadd_double_no_nopsub_pzero
+; CHECK-LABEL: define noundef nofpclass(pzero psub) float @fadd_double_no_nopsub_pzero
; CHECK-SAME: (float noundef nofpclass(pzero psub) [[ARG:%.*]]) #[[ATTR3]] {
; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
; CHECK-NEXT: ret float [[ADD]]
@@ -3196,7 +3196,7 @@ define float @fadd_double_no_nopsub_pzero(float noundef nofpclass(psub pzero) %a
define float @fadd_double_no_nonsub_nzero(float noundef nofpclass(nsub nzero) %arg) {
; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
-; CHECK-LABEL: define noundef nofpclass(nzero) float @fadd_double_no_nonsub_nzero
+; CHECK-LABEL: define noundef nofpclass(nzero nsub) float @fadd_double_no_nonsub_nzero
; CHECK-SAME: (float noundef nofpclass(nzero nsub) [[ARG:%.*]]) #[[ATTR3]] {
; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
; CHECK-NEXT: ret float [[ADD]]
@@ -3206,7 +3206,7 @@ define float @fadd_double_no_nonsub_nzero(float noundef nofpclass(nsub nzero) %a
}
define float @fadd_double_no_nopsub_pzero__ieee_daz(float noundef nofpclass(psub pzero) %arg) #2 {
-; CHECK-LABEL: define noundef nofpclass(pzero) float @fadd_double_no_nopsub_pzero__ieee_daz
+; CHECK-LABEL: define noundef nofpclass(pzero psub) float @fadd_double_no_nopsub_pzero__ieee_daz
; CHECK-SAME: (float noundef nofpclass(pzero psub) [[ARG:%.*]]) #[[ATTR11]] {
; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
; CHECK-NEXT: ret float [[ADD]]
@@ -3217,7 +3217,7 @@ define float @fadd_double_no_nopsub_pzero__ieee_daz(float noundef nofpclass(psub
define float @fadd_double_no_nopsub_pzero__ftz_daz(float noundef nofpclass(psub pzero) %arg) #0 {
; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn denormal_fpenv(preservesign) memory(none)
-; CHECK-LABEL: define noundef nofpclass(pzero) float @fadd_double_no_nopsub_pzero__ftz_daz
+; CHECK-LABEL: define noundef nofpclass(pzero psub) float @fadd_double_no_nopsub_pzero__ftz_daz
; CHECK-SAME: (float noundef nofpclass(pzero psub) [[ARG:%.*]]) #[[ATTR10]] {
; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
; CHECK-NEXT: ret float [[ADD]]
@@ -3227,7 +3227,7 @@ define float @fadd_double_no_nopsub_pzero__ftz_daz(float noundef nofpclass(psub
}
define float @fadd_double_no_nonsub_nzero__ieee_daz(float noundef nofpclass(nsub nzero) %arg) #2 {
-; CHECK-LABEL: define noundef nofpclass(nzero) float @fadd_double_no_nonsub_nzero__ieee_daz
+; CHECK-LABEL: define noundef nofpclass(nzero nsub) float @fadd_double_no_nonsub_nzero__ieee_daz
; CHECK-SAME: (float noundef nofpclass(nzero nsub) [[ARG:%.*]]) #[[ATTR11]] {
; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
; CHECK-NEXT: ret float [[ADD]]
@@ -3238,7 +3238,7 @@ define float @fadd_double_no_nonsub_nzero__ieee_daz(float noundef nofpclass(nsub
define float @fadd_double_no_nonsub_nzero__ftz_daz(float noundef nofpclass(nsub nzero) %arg) #0 {
; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn denormal_fpenv(preservesign) memory(none)
-; CHECK-LABEL: define noundef float @fadd_double_no_nonsub_nzero__ftz_daz
+; CHECK-LABEL: define noundef nofpclass(nzero nsub) float @fadd_double_no_nonsub_nzero__ftz_daz
; CHECK-SAME: (float noundef nofpclass(nzero nsub) [[ARG:%.*]]) #[[ATTR10]] {
; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
; CHECK-NEXT: ret float [[ADD]]
@@ -3248,7 +3248,7 @@ define float @fadd_double_no_nonsub_nzero__ftz_daz(float noundef nofpclass(nsub
}
define float @fadd_double_no_nopsub_pzero__ieee_dynamic(float noundef nofpclass(psub pzero) %arg) #9 {
-; CHECK-LABEL: define noundef float @fadd_double_no_nopsub_pzero__ieee_dynamic
+; CHECK-LABEL: define noundef nofpclass(psub) float @fadd_double_no_nopsub_pzero__ieee_dynamic
; CHECK-SAME: (float noundef nofpclass(pzero psub) [[ARG:%.*]]) #[[ATTR17:[0-9]+]] {
; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
; CHECK-NEXT: ret float [[ADD]]
@@ -3258,7 +3258,7 @@ define float @fadd_double_no_nopsub_pzero__ieee_dynamic(float noundef nofpclass(
}
define float @fadd_double_no_nonsub_nzero__ieee_dynamic(float noundef nofpclass(nsub nzero) %arg) #9 {
-; CHECK-LABEL: define noundef nofpclass(nzero) float @fadd_double_no_nonsub_nzero__ieee_dynamic
+; CHECK-LABEL: define noundef nofpclass(nzero nsub) float @fadd_double_no_nonsub_nzero__ieee_dynamic
; CHECK-SAME: (float noundef nofpclass(nzero nsub) [[ARG:%.*]]) #[[ATTR17]] {
; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
; CHECK-NEXT: ret float [[ADD]]
@@ -3739,7 +3739,7 @@ define float @fadd_double_no_zero__output_only_is_ftpz(float noundef nofpclass(z
}
define float @fadd_double_no_zero_or_nsub__output_only_is_ftpz(float noundef nofpclass(zero nsub) %arg) #4 {
-; CHECK-LABEL: define noundef nofpclass(nzero) float @fadd_double_no_zero_or_nsub__output_only_is_ftpz
+; CHECK-LABEL: define noundef nofpclass(nzero nsub) float @fadd_double_no_zero_or_nsub__output_only_is_ftpz
; CHECK-SAME: (float noundef nofpclass(zero nsub) [[ARG:%.*]]) #[[ATTR12]] {
; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
; CHECK-NEXT: ret float [[ADD]]
@@ -3749,7 +3749,7 @@ define float @fadd_double_no_zero_or_nsub__output_only_is_ftpz(float noundef nof
}
define float @fadd_double_no_zero_or_psub__output_only_is_ftpz(float noundef nofpclass(zero psub) %arg) #4 {
-; CHECK-LABEL: define noundef nofpclass(nzero) float @fadd_double_no_zero_or_psub__output_only_is_ftpz
+; CHECK-LABEL: define noundef nofpclass(nzero psub) float @fadd_double_no_zero_or_psub__output_only_is_ftpz
; CHECK-SAME: (float noundef nofpclass(zero psub) [[ARG:%.*]]) #[[ATTR12]] {
; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
; CHECK-NEXT: ret float [[ADD]]
@@ -3759,7 +3759,7 @@ define float @fadd_double_no_zero_or_psub__output_only_is_ftpz(float noundef nof
}
define float @fadd_double_no_zero_or_sub__output_only_is_ftpz(float noundef nofpclass(zero sub) %arg) #4 {
-; CHECK-LABEL: define noundef nofpclass(zero) float @fadd_double_no_zero_or_sub__output_only_is_ftpz
+; CHECK-LABEL: define noundef nofpclass(zero sub) float @fadd_double_no_zero_or_sub__output_only_is_ftpz
; CHECK-SAME: (float noundef nofpclass(zero sub) [[ARG:%.*]]) #[[ATTR12]] {
; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
; CHECK-NEXT: ret float [[ADD]]
@@ -3768,6 +3768,54 @@ define float @fadd_double_no_zero_or_sub__output_only_is_ftpz(float noundef nofp
ret float %add
}
+; Doubling never underflows. A subnormal result keeps the source sign.
+define float @fadd_double_no_nsub(float noundef nofpclass(nsub) %arg) {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
+; CHECK-LABEL: define noundef nofpclass(nsub) float @fadd_double_no_nsub
+; CHECK-SAME: (float noundef nofpclass(nsub) [[ARG:%.*]]) #[[ATTR3]] {
+; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
+; CHECK-NEXT: ret float [[ADD]]
+;
+ %add = fadd float %arg, %arg
+ ret float %add
+}
+
+; A dynamic output mode can flush either subnormal. Ruling both out also
+; rules out the zero it would produce.
+define float @fadd_double_no_pzero_or_sub__output_dynamic(float noundef nofpclass(pzero sub) %arg) #8 {
+; CHECK-LABEL: define noundef nofpclass(pzero sub) float @fadd_double_no_pzero_or_sub__output_dynamic
+; CHECK-SAME: (float noundef nofpclass(pzero sub) [[ARG:%.*]]) #[[ATTR18]] {
+; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
+; CHECK-NEXT: ret float [[ADD]]
+;
+ %add = fadd float %arg, %arg
+ ret float %add
+}
+
+; Negative test, double-double reports no precision, so doubling is not known
+; to leave the subnormal range.
+define ppc_fp128 @fadd_double_ppcf128_no_zero(ppc_fp128 noundef nofpclass(zero) %arg) {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
+; CHECK-LABEL: define noundef nofpclass(zero) ppc_fp128 @fadd_double_ppcf128_no_zero
+; CHECK-SAME: (ppc_fp128 noundef nofpclass(zero) [[ARG:%.*]]) #[[ATTR3]] {
+; CHECK-NEXT: [[ADD:%.*]] = fadd ppc_fp128 [[ARG]], [[ARG]]
+; CHECK-NEXT: ret ppc_fp128 [[ADD]]
+;
+ %add = fadd ppc_fp128 %arg, %arg
+ ret ppc_fp128 %add
+}
+
+; A preserve-sign output only reaches -0.0 through a negative subnormal.
+define float @fadd_double_no_nzero_or_nsub__output_ftz(float noundef nofpclass(nzero nsub) %arg) #1 {
+; CHECK-LABEL: define noundef nofpclass(nzero nsub) float @fadd_double_no_nzero_or_nsub__output_ftz
+; CHECK-SAME: (float noundef nofpclass(nzero nsub) [[ARG:%.*]]) #[[ATTR13]] {
+; CHECK-NEXT: [[ADD:%.*]] = fadd float [[ARG]], [[ARG]]
+; CHECK-NEXT: ret float [[ADD]]
+;
+ %add = fadd float %arg, %arg
+ ret float %add
+}
+
; Know there cannot be underflow, infer nofpclass(zero)
define half @known_positive_or_nan__fadd__known_positive_normal_or_inf(half nofpclass(ninf nnorm nsub nzero) %known.positive.or.nan, half nofpclass(ninf nnorm sub zero) %known.pnorm.or.pinf.or.nan) {
; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
diff --git a/llvm/test/Transforms/InstCombine/fptoui-of-fdiv.ll b/llvm/test/Transforms/InstCombine/fptoui-of-fdiv.ll
index 180b5915c05b1..a87d12ad2e898 100644
--- a/llvm/test/Transforms/InstCombine/fptoui-of-fdiv.ll
+++ b/llvm/test/Transforms/InstCombine/fptoui-of-fdiv.ll
@@ -110,10 +110,7 @@ define i32 @fractional_divisor(i32 %x) {
define i32 @zero_divisor(i32 %x) {
; CHECK-LABEL: define i32 @zero_divisor(
; CHECK-SAME: i32 [[X:%.*]]) {
-; CHECK-NEXT: [[A:%.*]] = uitofp i32 [[X]] to double
-; CHECK-NEXT: [[D:%.*]] = fdiv double [[A]], 0.000000e+00
-; CHECK-NEXT: [[R:%.*]] = fptoui double [[D]] to i32
-; CHECK-NEXT: ret i32 [[R]]
+; CHECK-NEXT: ret i32 0
;
%a = uitofp i32 %x to double
%d = fdiv double %a, 0.000000e+00
diff --git a/llvm/unittests/CodeGen/GlobalISel/KnownFPClassTest.cpp b/llvm/unittests/CodeGen/GlobalISel/KnownFPClassTest.cpp
index 9e3ed2f159b6f..92e190b3fc41e 100644
--- a/llvm/unittests/CodeGen/GlobalISel/KnownFPClassTest.cpp
+++ b/llvm/unittests/CodeGen/GlobalISel/KnownFPClassTest.cpp
@@ -1714,6 +1714,29 @@ TEST_F(AArch64GISelMITest, TestFPClassFMulAbsULEOne) {
EXPECT_EQ(std::nullopt, Known.SignBit);
}
+// G_FDIV by a constant of magnitude > 1 cannot grow the result. No infinity
+// or normal appears that the source does not have.
+TEST_F(AArch64GISelMITest, TestFPClassFDivByConstPow2) {
+ StringRef MIRString = R"(
+ %ptr:_(p0) = G_IMPLICIT_DEF
+ %val:_(s32) = G_LOAD %ptr(p0) :: (load (s32))
+ %finite:_(s32) = nnan ninf G_FABS %val
+ %c:_(s32) = G_FCONSTANT float 2.0
+ %fdiv:_(s32) = G_FDIV %finite, %c
+ %copy:_(s32) = COPY %fdiv
+)";
+ setUp(MIRString);
+ if (!TM)
+ GTEST_SKIP();
+ Register CopyReg = Copies[Copies.size() - 1];
+ MachineInstr *FinalCopy = MRI->getVRegDef(CopyReg);
+ Register SrcReg = FinalCopy->getOperand(1).getReg();
+ GISelValueTracking Info(*MF);
+ KnownFPClass Known = Info.computeKnownFPClass(SrcReg);
+ EXPECT_EQ(fcZero | fcSubnormal | fcPosNormal, Known.KnownFPClasses);
+ EXPECT_EQ(false, Known.SignBit);
+}
+
// G_FMA with A == B (and A guaranteed not-undef): the multiply part is a
// square, so the result is known non-negative (never fcNegative).
TEST_F(AArch64GISelMITest, TestFPClassFMASelfSquare) {
>From 05cb5ede987b5d69696bc9aee9dd31a9d69a9124 Mon Sep 17 00:00:00 2001
From: MaxGraey <maxgraey at gmail.com>
Date: Sun, 9 Aug 2026 16:47:52 +0300
Subject: [PATCH 2/5] minor fixes
---
llvm/lib/Support/KnownFPClass.cpp | 9 +++++----
1 file changed, 5 insertions(+), 4 deletions(-)
diff --git a/llvm/lib/Support/KnownFPClass.cpp b/llvm/lib/Support/KnownFPClass.cpp
index f9522fadfcf02..c1cea0282ebf3 100644
--- a/llvm/lib/Support/KnownFPClass.cpp
+++ b/llvm/lib/Support/KnownFPClass.cpp
@@ -401,8 +401,7 @@ KnownFPClass KnownFPClass::fadd_self(const KnownFPClass &KnownSrc,
(fcZero | fcSubnormal))
return Known;
- // x + x scales by 2, down to 1 when one read flushes and leaves x + 0 = x.
- Known.propagateExpRange(KnownSrc, 0, 1, /*NegativeScale=*/false, Sem, Mode);
+ Known.propagateExpRange(KnownSrc, 1, 1, /*NegativeScale=*/false, Sem, Mode);
return Known;
}
@@ -455,7 +454,8 @@ KnownFPClass KnownFPClass::fmul(const KnownFPClass &KnownLHS,
const int Exp = ilogb(ConstRHS);
const bool IsPow2 = ConstRHS.getExactLog2Abs() != INT_MIN;
- // |C| in [2^Exp, 2^(Exp+1)), exactly 2^Exp for a power of two.
+ // |C| in [2^Exp, 2^Exp] exactly when C is power of two
+ // |C| in [2^Exp, 2^(Exp+1)) otherwise
Known.propagateExpRange(KnownLHS, Exp, IsPow2 ? Exp : Exp + 1,
ConstRHS.isNegative(), ConstRHS.getSemantics(), Mode);
return Known;
@@ -501,7 +501,8 @@ KnownFPClass KnownFPClass::fdiv(const KnownFPClass &KnownLHS,
const int Exp = ilogb(ConstRHS);
const bool IsPow2 = ConstRHS.getExactLog2Abs() != INT_MIN;
- // |1 / C| in (2^(-Exp-1), 2^-Exp], exactly 2^-Exp for a power of two.
+ // |1 / C| in [2^-Exp, 2^-Exp] exactly when C is power of two
+ // |1 / C| in (2^(-Exp-1), 2^-Exp] otherwise
Known.propagateExpRange(KnownLHS, IsPow2 ? -Exp : -Exp - 1, -Exp,
ConstRHS.isNegative(), ConstRHS.getSemantics(), Mode);
return Known;
>From b4577ffa3ab6a862568b0d720d300345db71a72b Mon Sep 17 00:00:00 2001
From: MaxGraey <maxgraey at gmail.com>
Date: Sun, 9 Aug 2026 17:48:25 +0300
Subject: [PATCH 3/5] fix GISel test
---
llvm/unittests/CodeGen/GlobalISel/KnownFPClassTest.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/unittests/CodeGen/GlobalISel/KnownFPClassTest.cpp b/llvm/unittests/CodeGen/GlobalISel/KnownFPClassTest.cpp
index 92e190b3fc41e..d21b02d3f4cc0 100644
--- a/llvm/unittests/CodeGen/GlobalISel/KnownFPClassTest.cpp
+++ b/llvm/unittests/CodeGen/GlobalISel/KnownFPClassTest.cpp
@@ -1733,7 +1733,7 @@ TEST_F(AArch64GISelMITest, TestFPClassFDivByConstPow2) {
Register SrcReg = FinalCopy->getOperand(1).getReg();
GISelValueTracking Info(*MF);
KnownFPClass Known = Info.computeKnownFPClass(SrcReg);
- EXPECT_EQ(fcZero | fcSubnormal | fcPosNormal, Known.KnownFPClasses);
+ EXPECT_EQ(fcPosZero | fcPosSubnormal | fcPosNormal, Known.KnownFPClasses);
EXPECT_EQ(false, Known.SignBit);
}
>From 5694279a83104f1f8b89ff4ecb6c461b8d0a83bc Mon Sep 17 00:00:00 2001
From: MaxGraey <maxgraey at gmail.com>
Date: Sun, 9 Aug 2026 17:53:13 +0300
Subject: [PATCH 4/5] rename NegativeScale to IsNegative
---
llvm/include/llvm/Support/KnownFPClass.h | 4 ++--
llvm/lib/Support/KnownFPClass.cpp | 8 ++++----
2 files changed, 6 insertions(+), 6 deletions(-)
diff --git a/llvm/include/llvm/Support/KnownFPClass.h b/llvm/include/llvm/Support/KnownFPClass.h
index ed4af9b18af29..3ed4548e39d2e 100644
--- a/llvm/include/llvm/Support/KnownFPClass.h
+++ b/llvm/include/llvm/Support/KnownFPClass.h
@@ -412,9 +412,9 @@ struct KnownFPClass {
/// Refine the known classes of Src scaled by a factor with magnitude in
/// [2^LoExp, 2^HiExp], as ldexp and a constant multiply or divide do. A
/// scale that cannot grow rules out infinities and normals, one that cannot
- /// shrink rules out subnormals and zeroes. NegativeScale flips the sign.
+ /// shrink rules out subnormals and zeroes. IsNegative flips the sign.
LLVM_ABI void propagateExpRange(const KnownFPClass &Src, int LoExp, int HiExp,
- bool NegativeScale, const fltSemantics &Sem,
+ bool IsNegative, const fltSemantics &Sem,
DenormalMode Mode);
/// Propagate knowledge from a source value that could be a denormal or
diff --git a/llvm/lib/Support/KnownFPClass.cpp b/llvm/lib/Support/KnownFPClass.cpp
index c1cea0282ebf3..e8fd8cacc9d61 100644
--- a/llvm/lib/Support/KnownFPClass.cpp
+++ b/llvm/lib/Support/KnownFPClass.cpp
@@ -68,7 +68,7 @@ bool KnownFPClass::isKnownNeverLogicalPosZero(DenormalMode Mode) const {
}
void KnownFPClass::propagateExpRange(const KnownFPClass &Src, int LoExp,
- int HiExp, bool NegativeScale,
+ int HiExp, bool IsNegative,
const fltSemantics &Sem,
DenormalMode Mode) {
FPClassTest Result = fcNone;
@@ -79,7 +79,7 @@ void KnownFPClass::propagateExpRange(const KnownFPClass &Src, int LoExp,
if (!Src.isKnownNeverLogicalNegZero(Mode))
KnownFPClasses |= fcNegZero;
- if (NegativeScale)
+ if (IsNegative)
KnownFPClasses = llvm::fneg(KnownFPClasses);
// A scale that cannot grow keeps |result| <= |source|.
@@ -401,7 +401,7 @@ KnownFPClass KnownFPClass::fadd_self(const KnownFPClass &KnownSrc,
(fcZero | fcSubnormal))
return Known;
- Known.propagateExpRange(KnownSrc, 1, 1, /*NegativeScale=*/false, Sem, Mode);
+ Known.propagateExpRange(KnownSrc, 1, 1, /*IsNegative=*/false, Sem, Mode);
return Known;
}
@@ -882,7 +882,7 @@ KnownFPClass KnownFPClass::ldexp(const KnownFPClass &KnownSrc,
Known.propagateExpRange(KnownSrc, SaturatedIntExp(ConstantRangeExpMin),
SaturatedIntExp(ConstantRangeExpMax),
- /*NegativeScale=*/false, Flt, Mode);
+ /*IsNegative=*/false, Flt, Mode);
return Known;
}
>From 374f74c9ab4b367c265713af5fd773988a363d00 Mon Sep 17 00:00:00 2001
From: MaxGraey <maxgraey at gmail.com>
Date: Mon, 10 Aug 2026 14:42:39 +0300
Subject: [PATCH 5/5] refactor
---
llvm/lib/Support/KnownFPClass.cpp | 5 +++--
1 file changed, 3 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Support/KnownFPClass.cpp b/llvm/lib/Support/KnownFPClass.cpp
index e8fd8cacc9d61..804a37f68d775 100644
--- a/llvm/lib/Support/KnownFPClass.cpp
+++ b/llvm/lib/Support/KnownFPClass.cpp
@@ -90,9 +90,10 @@ void KnownFPClass::propagateExpRange(const KnownFPClass &Src, int LoExp,
// 2^MantissaBits lifts even the smallest value 2^(emin - MantissaBits) into
// the normal range. Double-double reports no precision and has no such
// threshold, only its source can rule a subnormal out.
- const unsigned Precision = APFloat::semanticsPrecision(Sem);
+ const int MantissaBits =
+ static_cast<int>(APFloat::semanticsPrecision(Sem)) - 1;
const bool LiftsOutOfSubnormals =
- Precision != 0 && LoExp >= static_cast<int>(Precision) - 1;
+ MantissaBits >= 0 && LoExp >= MantissaBits;
Result |=
LiftsOutOfSubnormals ? fcSubnormal : ~KnownFPClasses & fcSubnormal;
More information about the llvm-commits
mailing list