From 9be777d5b39852cf3c0b2538fd5f712922672caa Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Fri, 1 Dec 2023 18:00:13 +0900 Subject: [PATCH 1/4] Reapply "InstCombine: Introduce SimplifyDemandedUseFPClass"" This reverts commit ef388334ee5a3584255b9ef5b3fefdb244fa3fd7. The referenced issue violates the spec for finite-only math only by using a return value for a constant infinity. --- clang/test/Headers/__clang_hip_math.hip | 56 +++-- llvm/include/llvm/Analysis/ValueTracking.h | 4 + .../InstCombine/InstCombineInternal.h | 9 + .../InstCombineSimplifyDemanded.cpp | 140 +++++++++++- .../InstCombine/InstructionCombining.cpp | 18 +- .../InstCombine/simplify-demanded-fpclass.ll | 203 +++++++----------- 6 files changed, 286 insertions(+), 144 deletions(-) diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index c0f4a06acbb8e..9e15aec94dc28 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -2557,33 +2557,65 @@ extern "C" __device__ double test_nan(const char *tag) { return nan(tag); } -// CHECK-LABEL: @test_nanf_emptystr( -// CHECK-NEXT: entry: -// CHECK-NEXT: ret float 0x7FF8000000000000 +// DEFAULT-LABEL: @test_nanf_emptystr( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: ret float 0x7FF8000000000000 +// +// FINITEONLY-LABEL: @test_nanf_emptystr( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: ret float poison +// +// APPROX-LABEL: @test_nanf_emptystr( +// APPROX-NEXT: entry: +// APPROX-NEXT: ret float 0x7FF8000000000000 // extern "C" __device__ float test_nanf_emptystr() { return nanf(""); } -// CHECK-LABEL: @test_nan_emptystr( -// CHECK-NEXT: entry: -// CHECK-NEXT: ret double 0x7FF8000000000000 +// DEFAULT-LABEL: @test_nan_emptystr( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: ret double 0x7FF8000000000000 +// +// FINITEONLY-LABEL: @test_nan_emptystr( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: ret double poison +// +// APPROX-LABEL: @test_nan_emptystr( +// APPROX-NEXT: entry: +// APPROX-NEXT: ret double 0x7FF8000000000000 // extern "C" __device__ double test_nan_emptystr() { return nan(""); } -// CHECK-LABEL: @test_nanf_fill( -// CHECK-NEXT: entry: -// CHECK-NEXT: ret float 0x7FF8000000000000 +// DEFAULT-LABEL: @test_nanf_fill( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: ret float 0x7FF8000000000000 +// +// FINITEONLY-LABEL: @test_nanf_fill( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: ret float poison +// +// APPROX-LABEL: @test_nanf_fill( +// APPROX-NEXT: entry: +// APPROX-NEXT: ret float 0x7FF8000000000000 // extern "C" __device__ float test_nanf_fill() { return nanf("0x456"); } -// CHECK-LABEL: @test_nan_fill( -// CHECK-NEXT: entry: -// CHECK-NEXT: ret double 0x7FF8000000000000 +// DEFAULT-LABEL: @test_nan_fill( +// DEFAULT-NEXT: entry: +// DEFAULT-NEXT: ret double 0x7FF8000000000000 +// +// FINITEONLY-LABEL: @test_nan_fill( +// FINITEONLY-NEXT: entry: +// FINITEONLY-NEXT: ret double poison +// +// APPROX-LABEL: @test_nan_fill( +// APPROX-NEXT: entry: +// APPROX-NEXT: ret double 0x7FF8000000000000 // extern "C" __device__ double test_nan_fill() { return nan("0x123"); diff --git a/llvm/include/llvm/Analysis/ValueTracking.h b/llvm/include/llvm/Analysis/ValueTracking.h index 82c87edd6297c..f9ce679bc7426 100644 --- a/llvm/include/llvm/Analysis/ValueTracking.h +++ b/llvm/include/llvm/Analysis/ValueTracking.h @@ -243,6 +243,10 @@ struct KnownFPClass { /// definitely set or false if the sign bit is definitely unset. std::optional SignBit; + bool operator==(KnownFPClass Other) const { + return KnownFPClasses == Other.KnownFPClasses && SignBit == Other.SignBit; + } + /// Return true if it's known this can never be one of the mask entries. bool isKnownNever(FPClassTest Mask) const { return (KnownFPClasses & Mask) == fcNone; diff --git a/llvm/lib/Transforms/InstCombine/InstCombineInternal.h b/llvm/lib/Transforms/InstCombine/InstCombineInternal.h index 0bbb22be71569..9a66fb8f456f9 100644 --- a/llvm/lib/Transforms/InstCombine/InstCombineInternal.h +++ b/llvm/lib/Transforms/InstCombine/InstCombineInternal.h @@ -551,6 +551,15 @@ class LLVM_LIBRARY_VISIBILITY InstCombinerImpl final APInt &UndefElts, unsigned Depth = 0, bool AllowMultipleUsers = false) override; + /// Attempts to replace V with a simpler value based on the demanded + /// floating-point classes + Value *SimplifyDemandedUseFPClass(Value *V, FPClassTest DemandedMask, + KnownFPClass &Known, unsigned Depth, + Instruction *CxtI); + bool SimplifyDemandedFPClass(Instruction *I, unsigned Op, + FPClassTest DemandedMask, KnownFPClass &Known, + unsigned Depth = 0); + /// Canonicalize the position of binops relative to shufflevector. Instruction *foldVectorBinop(BinaryOperator &Inst); Instruction *foldVectorSelect(SelectInst &Sel); diff --git a/llvm/lib/Transforms/InstCombine/InstCombineSimplifyDemanded.cpp b/llvm/lib/Transforms/InstCombine/InstCombineSimplifyDemanded.cpp index 4d19bd12d8f6f..418155e90a2a7 100644 --- a/llvm/lib/Transforms/InstCombine/InstCombineSimplifyDemanded.cpp +++ b/llvm/lib/Transforms/InstCombine/InstCombineSimplifyDemanded.cpp @@ -466,7 +466,8 @@ Value *InstCombinerImpl::SimplifyDemandedUseBits(Value *V, APInt DemandedMask, if (InputKnown.isNonNegative() || DemandedMask.getActiveBits() <= SrcBitWidth) { // Convert to ZExt cast. - CastInst *NewCast = new ZExtInst(I->getOperand(0), VTy, I->getName()); + CastInst *NewCast = new ZExtInst(I->getOperand(0), VTy); + NewCast->takeName(I); return InsertNewInstWith(NewCast, I->getIterator()); } @@ -774,6 +775,7 @@ Value *InstCombinerImpl::SimplifyDemandedUseBits(Value *V, APInt DemandedMask, BinaryOperator *LShr = BinaryOperator::CreateLShr(I->getOperand(0), I->getOperand(1)); LShr->setIsExact(cast(I)->isExact()); + LShr->takeName(I); return InsertNewInstWith(LShr, I->getIterator()); } else if (Known.One[BitWidth-ShiftAmt-1]) { // New bits are known one. Known.One |= HighBits; @@ -1849,3 +1851,139 @@ Value *InstCombinerImpl::SimplifyDemandedVectorElts(Value *V, return MadeChange ? I : nullptr; } + +/// For floating-point classes that resolve to a single bit pattern, return that +/// value. +static Constant *getFPClassConstant(Type *Ty, FPClassTest Mask) { + switch (Mask) { + case fcPosZero: + return ConstantFP::getZero(Ty); + case fcNegZero: + return ConstantFP::getZero(Ty, true); + case fcPosInf: + return ConstantFP::getInfinity(Ty); + case fcNegInf: + return ConstantFP::getInfinity(Ty, true); + case fcNone: + return PoisonValue::get(Ty); + default: + return nullptr; + } +} + +Value *InstCombinerImpl::SimplifyDemandedUseFPClass( + Value *V, const FPClassTest DemandedMask, KnownFPClass &Known, + unsigned Depth, Instruction *CxtI) { + assert(Depth <= MaxAnalysisRecursionDepth && "Limit Search Depth"); + Type *VTy = V->getType(); + + assert(Known == KnownFPClass() && "expected uninitialized state"); + + if (DemandedMask == fcNone) + return isa(V) ? nullptr : PoisonValue::get(VTy); + + if (Depth == MaxAnalysisRecursionDepth) + return nullptr; + + Instruction *I = dyn_cast(V); + if (!I) { + // Handle constants and arguments + Known = computeKnownFPClass(V, fcAllFlags, CxtI, Depth + 1); + Value *FoldedToConst = + getFPClassConstant(VTy, DemandedMask & Known.KnownFPClasses); + return FoldedToConst == V ? nullptr : FoldedToConst; + } + + if (!I->hasOneUse()) + return nullptr; + + // TODO: Should account for nofpclass/FastMathFlags on current instruction + switch (I->getOpcode()) { + case Instruction::FNeg: { + if (SimplifyDemandedFPClass(I, 0, llvm::fneg(DemandedMask), Known, + Depth + 1)) + return I; + Known.fneg(); + break; + } + case Instruction::Call: { + CallInst *CI = cast(I); + switch (CI->getIntrinsicID()) { + case Intrinsic::fabs: + if (SimplifyDemandedFPClass(I, 0, llvm::inverse_fabs(DemandedMask), Known, + Depth + 1)) + return I; + Known.fabs(); + break; + case Intrinsic::arithmetic_fence: + if (SimplifyDemandedFPClass(I, 0, DemandedMask, Known, Depth + 1)) + return I; + break; + case Intrinsic::copysign: { + // Flip on more potentially demanded classes + const FPClassTest DemandedMaskAnySign = llvm::unknown_sign(DemandedMask); + if (SimplifyDemandedFPClass(I, 0, DemandedMaskAnySign, Known, Depth + 1)) + return I; + + if ((DemandedMask & fcPositive) == fcNone) { + // Roundabout way of replacing with fneg(fabs) + I->setOperand(1, ConstantFP::get(VTy, -1.0)); + return I; + } + + if ((DemandedMask & fcNegative) == fcNone) { + // Roundabout way of replacing with fabs + I->setOperand(1, ConstantFP::getZero(VTy)); + return I; + } + + KnownFPClass KnownSign = + computeKnownFPClass(I->getOperand(1), fcAllFlags, CxtI, Depth + 1); + Known.copysign(KnownSign); + break; + } + default: + Known = computeKnownFPClass(I, ~DemandedMask, CxtI, Depth + 1); + break; + } + + break; + } + case Instruction::Select: { + KnownFPClass KnownLHS, KnownRHS; + if (SimplifyDemandedFPClass(I, 2, DemandedMask, KnownRHS, Depth + 1) || + SimplifyDemandedFPClass(I, 1, DemandedMask, KnownLHS, Depth + 1)) + return I; + + if (KnownLHS.isKnownNever(DemandedMask)) + return I->getOperand(2); + if (KnownRHS.isKnownNever(DemandedMask)) + return I->getOperand(1); + + // TODO: Recognize clamping patterns + Known = KnownLHS | KnownRHS; + break; + } + default: + Known = computeKnownFPClass(I, ~DemandedMask, CxtI, Depth + 1); + break; + } + + return getFPClassConstant(VTy, DemandedMask & Known.KnownFPClasses); +} + +bool InstCombinerImpl::SimplifyDemandedFPClass(Instruction *I, unsigned OpNo, + FPClassTest DemandedMask, + KnownFPClass &Known, + unsigned Depth) { + Use &U = I->getOperandUse(OpNo); + Value *NewVal = + SimplifyDemandedUseFPClass(U.get(), DemandedMask, Known, Depth, I); + if (!NewVal) + return false; + if (Instruction *OpInst = dyn_cast(U)) + salvageDebugInfo(*OpInst); + + replaceUse(U, NewVal); + return true; +} diff --git a/llvm/lib/Transforms/InstCombine/InstructionCombining.cpp b/llvm/lib/Transforms/InstCombine/InstructionCombining.cpp index 26fdef672506a..4f2f4edea6e09 100644 --- a/llvm/lib/Transforms/InstCombine/InstructionCombining.cpp +++ b/llvm/lib/Transforms/InstCombine/InstructionCombining.cpp @@ -2904,8 +2904,22 @@ Instruction *InstCombinerImpl::visitFree(CallInst &FI, Value *Op) { } Instruction *InstCombinerImpl::visitReturnInst(ReturnInst &RI) { - // Nothing for now. - return nullptr; + Value *RetVal = RI.getReturnValue(); + if (!RetVal || !AttributeFuncs::isNoFPClassCompatibleType(RetVal->getType())) + return nullptr; + + Function *F = RI.getFunction(); + FPClassTest ReturnClass = F->getAttributes().getRetNoFPClass(); + if (ReturnClass == fcNone) + return nullptr; + + KnownFPClass KnownClass; + Value *Simplified = + SimplifyDemandedUseFPClass(RetVal, ~ReturnClass, KnownClass, 0, &RI); + if (!Simplified) + return nullptr; + + return ReturnInst::Create(RI.getContext(), Simplified); } // WARNING: keep in sync with SimplifyCFGOpt::simplifyUnreachable()! diff --git a/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass.ll b/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass.ll index 9817b6e13ca8a..4f9396add2370 100644 --- a/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass.ll +++ b/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass.ll @@ -42,7 +42,7 @@ define nofpclass(inf) float @ret_nofpclass_inf_undef() { define nofpclass(all) float @ret_nofpclass_all_var(float %arg) { ; CHECK-LABEL: define nofpclass(all) float @ret_nofpclass_all_var ; CHECK-SAME: (float [[ARG:%.*]]) { -; CHECK-NEXT: ret float [[ARG]] +; CHECK-NEXT: ret float poison ; ret float %arg } @@ -51,7 +51,7 @@ define nofpclass(all) float @ret_nofpclass_all_var(float %arg) { define nofpclass(all) <2 x float> @ret_nofpclass_all_var_vector(<2 x float> %arg) { ; CHECK-LABEL: define nofpclass(all) <2 x float> @ret_nofpclass_all_var_vector ; CHECK-SAME: (<2 x float> [[ARG:%.*]]) { -; CHECK-NEXT: ret <2 x float> [[ARG]] +; CHECK-NEXT: ret <2 x float> poison ; ret <2 x float> %arg } @@ -65,14 +65,14 @@ define nofpclass(inf) float @ret_nofpclass_inf__0() { define nofpclass(inf) float @ret_nofpclass_inf__pinf() { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__pinf() { -; CHECK-NEXT: ret float 0x7FF0000000000000 +; CHECK-NEXT: ret float poison ; ret float 0x7FF0000000000000 } define nofpclass(pinf) float @ret_nofpclass_pinf__pinf() { ; CHECK-LABEL: define nofpclass(pinf) float @ret_nofpclass_pinf__pinf() { -; CHECK-NEXT: ret float 0x7FF0000000000000 +; CHECK-NEXT: ret float poison ; ret float 0x7FF0000000000000 } @@ -86,7 +86,7 @@ define nofpclass(pinf) float @ret_nofpclass_pinf__ninf() { define nofpclass(inf) float @ret_nofpclass_inf__ninf() { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__ninf() { -; CHECK-NEXT: ret float 0xFFF0000000000000 +; CHECK-NEXT: ret float poison ; ret float 0xFFF0000000000000 } @@ -106,8 +106,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__select_nofpclass_inf_lhs(i1 %con define nofpclass(inf) float @ret_nofpclass_inf__select_nofpclass_arg_only_inf_lhs(i1 %cond, float nofpclass(nan norm zero sub) %x, float %y) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__select_nofpclass_arg_only_inf_lhs ; CHECK-SAME: (i1 [[COND:%.*]], float nofpclass(nan zero sub norm) [[X:%.*]], float [[Y:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float [[Y]] -; CHECK-NEXT: ret float [[SELECT]] +; CHECK-NEXT: ret float [[Y]] ; %select = select i1 %cond, float %x, float %y ret float %select @@ -117,8 +116,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__select_nofpclass_arg_only_inf_lh define nofpclass(inf) float @ret_nofpclass_inf__select_nofpclass_arg_only_inf_rhs(i1 %cond, float %x, float nofpclass(nan norm zero sub) %y) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__select_nofpclass_arg_only_inf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]], float nofpclass(nan zero sub norm) [[Y:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float [[Y]] -; CHECK-NEXT: ret float [[SELECT]] +; CHECK-NEXT: ret float [[X]] ; %select = select i1 %cond, float %x, float %y ret float %select @@ -128,8 +126,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__select_nofpclass_arg_only_inf_rh define nofpclass(inf) [3 x [2 x float]] @ret_float_array(i1 %cond, [3 x [2 x float]] nofpclass(nan norm zero sub) %x, [3 x [2 x float]] %y) { ; CHECK-LABEL: define nofpclass(inf) [3 x [2 x float]] @ret_float_array ; CHECK-SAME: (i1 [[COND:%.*]], [3 x [2 x float]] nofpclass(nan zero sub norm) [[X:%.*]], [3 x [2 x float]] [[Y:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], [3 x [2 x float]] [[X]], [3 x [2 x float]] [[Y]] -; CHECK-NEXT: ret [3 x [2 x float]] [[SELECT]] +; CHECK-NEXT: ret [3 x [2 x float]] [[Y]] ; %select = select i1 %cond, [3 x [2 x float]] %x, [3 x [2 x float]] %y ret [3 x [2 x float ]] %select @@ -139,8 +136,7 @@ define nofpclass(inf) [3 x [2 x float]] @ret_float_array(i1 %cond, [3 x [2 x flo define nofpclass(inf) float @ret_nofpclass_inf__select_pinf_lhs(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__select_pinf_lhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float 0x7FF0000000000000, float [[X]] -; CHECK-NEXT: ret float [[SELECT]] +; CHECK-NEXT: ret float [[X]] ; %select = select i1 %cond, float 0x7FF0000000000000, float %x ret float %select @@ -150,8 +146,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__select_pinf_lhs(i1 %cond, float define nofpclass(inf) float @ret_nofpclass_inf__select_pinf_rhs(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: ret float [[SELECT]] +; CHECK-NEXT: ret float [[X]] ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 ret float %select @@ -161,8 +156,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__select_pinf_rhs(i1 %cond, float define nofpclass(inf) float @ret_nofpclass_inf__select_pinf_or_ninf(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__select_pinf_or_ninf ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float 0x7FF0000000000000, float 0xFFF0000000000000 -; CHECK-NEXT: ret float [[SELECT]] +; CHECK-NEXT: ret float poison ; %select = select i1 %cond, float 0x7FF0000000000000, float 0xFFF0000000000000 ret float %select @@ -172,8 +166,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__select_pinf_or_ninf(i1 %cond, fl define nofpclass(inf) float @ret_nofpclass_inf__select_ninf_or_pinf(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__select_ninf_or_pinf ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float 0xFFF0000000000000, float 0x7FF0000000000000 -; CHECK-NEXT: ret float [[SELECT]] +; CHECK-NEXT: ret float poison ; %select = select i1 %cond, float 0xFFF0000000000000, float 0x7FF0000000000000 ret float %select @@ -183,8 +176,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__select_ninf_or_pinf(i1 %cond, fl define nofpclass(ninf) float @ret_nofpclass_ninf__select_ninf_or_pinf(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(ninf) float @ret_nofpclass_ninf__select_ninf_or_pinf ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float 0xFFF0000000000000, float 0x7FF0000000000000 -; CHECK-NEXT: ret float [[SELECT]] +; CHECK-NEXT: ret float 0x7FF0000000000000 ; %select = select i1 %cond, float 0xFFF0000000000000, float 0x7FF0000000000000 ret float %select @@ -194,8 +186,7 @@ define nofpclass(ninf) float @ret_nofpclass_ninf__select_ninf_or_pinf(i1 %cond, define nofpclass(pinf) float @ret_nofpclass_pinf__select_ninf_or_pinf(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(pinf) float @ret_nofpclass_pinf__select_ninf_or_pinf ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float 0xFFF0000000000000, float 0x7FF0000000000000 -; CHECK-NEXT: ret float [[SELECT]] +; CHECK-NEXT: ret float 0xFFF0000000000000 ; %select = select i1 %cond, float 0xFFF0000000000000, float 0x7FF0000000000000 ret float %select @@ -205,8 +196,7 @@ define nofpclass(pinf) float @ret_nofpclass_pinf__select_ninf_or_pinf(i1 %cond, define nofpclass(zero) float @ret_nofpclass_zero__select_pzero_or_nzero(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(zero) float @ret_nofpclass_zero__select_pzero_or_nzero ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float 0.000000e+00, float -0.000000e+00 -; CHECK-NEXT: ret float [[SELECT]] +; CHECK-NEXT: ret float poison ; %select = select i1 %cond, float 0.0, float -0.0 ret float %select @@ -216,8 +206,7 @@ define nofpclass(zero) float @ret_nofpclass_zero__select_pzero_or_nzero(i1 %cond define nofpclass(nzero) float @ret_nofpclass_nzero__select_pzero_or_nzero(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(nzero) float @ret_nofpclass_nzero__select_pzero_or_nzero ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float 0.000000e+00, float -0.000000e+00 -; CHECK-NEXT: ret float [[SELECT]] +; CHECK-NEXT: ret float 0.000000e+00 ; %select = select i1 %cond, float 0.0, float -0.0 ret float %select @@ -227,8 +216,7 @@ define nofpclass(nzero) float @ret_nofpclass_nzero__select_pzero_or_nzero(i1 %co define nofpclass(pzero) float @ret_nofpclass_pzero__select_pzero_or_nzero(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(pzero) float @ret_nofpclass_pzero__select_pzero_or_nzero ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float 0.000000e+00, float -0.000000e+00 -; CHECK-NEXT: ret float [[SELECT]] +; CHECK-NEXT: ret float -0.000000e+00 ; %select = select i1 %cond, float 0.0, float -0.0 ret float %select @@ -238,8 +226,7 @@ define nofpclass(pzero) float @ret_nofpclass_pzero__select_pzero_or_nzero(i1 %co define nofpclass(inf) <2 x float> @ret_nofpclass_inf__select_pinf_lhs_vector(<2 x i1> %cond, <2 x float> %x) { ; CHECK-LABEL: define nofpclass(inf) <2 x float> @ret_nofpclass_inf__select_pinf_lhs_vector ; CHECK-SAME: (<2 x i1> [[COND:%.*]], <2 x float> [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select <2 x i1> [[COND]], <2 x float> , <2 x float> [[X]] -; CHECK-NEXT: ret <2 x float> [[SELECT]] +; CHECK-NEXT: ret <2 x float> [[X]] ; %select = select <2 x i1> %cond, <2 x float> , <2 x float> %x ret <2 x float> %select @@ -249,8 +236,7 @@ define nofpclass(inf) <2 x float> @ret_nofpclass_inf__select_pinf_lhs_vector(<2 define nofpclass(inf) <2 x float> @ret_nofpclass_inf__select_pinf_lhs_vector_undef(<2 x i1> %cond, <2 x float> %x) { ; CHECK-LABEL: define nofpclass(inf) <2 x float> @ret_nofpclass_inf__select_pinf_lhs_vector_undef ; CHECK-SAME: (<2 x i1> [[COND:%.*]], <2 x float> [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select <2 x i1> [[COND]], <2 x float> , <2 x float> [[X]] -; CHECK-NEXT: ret <2 x float> [[SELECT]] +; CHECK-NEXT: ret <2 x float> [[X]] ; %select = select <2 x i1> %cond, <2 x float> , <2 x float> %x ret <2 x float> %select @@ -260,8 +246,7 @@ define nofpclass(inf) <2 x float> @ret_nofpclass_inf__select_pinf_lhs_vector_und define nofpclass(inf) <2 x float> @ret_nofpclass_inf__select_mixed_inf_lhs_vector(<2 x i1> %cond, <2 x float> %x) { ; CHECK-LABEL: define nofpclass(inf) <2 x float> @ret_nofpclass_inf__select_mixed_inf_lhs_vector ; CHECK-SAME: (<2 x i1> [[COND:%.*]], <2 x float> [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select <2 x i1> [[COND]], <2 x float> , <2 x float> [[X]] -; CHECK-NEXT: ret <2 x float> [[SELECT]] +; CHECK-NEXT: ret <2 x float> [[X]] ; %select = select <2 x i1> %cond, <2 x float> , <2 x float> %x ret <2 x float> %select @@ -327,8 +312,7 @@ define nofpclass(nan) float @ret_nofpclass_nan__select_pinf_rhs(i1 %cond, float define nofpclass(inf nan) float @ret_nofpclass_inf_nan__select_chain_inf_nan_0(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(nan inf) float @ret_nofpclass_inf_nan__select_chain_inf_nan_0 ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT1:%.*]] = select i1 [[COND]], float 0x7FF0000000000000, float [[X]] -; CHECK-NEXT: ret float [[SELECT1]] +; CHECK-NEXT: ret float [[X]] ; %select0 = select i1 %cond, float 0x7FF8000000000000, float %x %select1 = select i1 %cond, float 0x7FF0000000000000, float %select0 @@ -338,8 +322,7 @@ define nofpclass(inf nan) float @ret_nofpclass_inf_nan__select_chain_inf_nan_0(i define nofpclass(inf nan) float @ret_nofpclass_inf_nan__select_chain_inf_nan_1(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(nan inf) float @ret_nofpclass_inf_nan__select_chain_inf_nan_1 ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT1:%.*]] = select i1 [[COND]], float 0x7FF0000000000000, float 0x7FF8000000000000 -; CHECK-NEXT: ret float [[SELECT1]] +; CHECK-NEXT: ret float poison ; %select0 = select i1 %cond, float %x, float 0x7FF8000000000000 %select1 = select i1 %cond, float 0x7FF0000000000000, float %select0 @@ -360,8 +343,7 @@ define nofpclass(nan) float @ret_nofpclass_nan__select_chain_inf_nan(i1 %cond, f define nofpclass(inf) float @ret_nofpclass_inf__select_chain_inf_nan_0(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__select_chain_inf_nan_0 ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT1:%.*]] = select i1 [[COND]], float 0x7FF0000000000000, float [[X]] -; CHECK-NEXT: ret float [[SELECT1]] +; CHECK-NEXT: ret float [[X]] ; %select0 = select i1 %cond, float 0x7FF8000000000000, float %x %select1 = select i1 %cond, float 0x7FF0000000000000, float %select0 @@ -371,8 +353,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__select_chain_inf_nan_0(i1 %cond, define nofpclass(inf) float @ret_nofpclass_inf__select_chain_inf_nan_1(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__select_chain_inf_nan_1 ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT1:%.*]] = select i1 [[COND]], float 0x7FF8000000000000, float 0x7FF0000000000000 -; CHECK-NEXT: ret float [[SELECT1]] +; CHECK-NEXT: ret float 0x7FF8000000000000 ; %select0 = select i1 %cond, float 0x7FF8000000000000, float %x %select1 = select i1 %cond, float %select0, float 0x7FF0000000000000 @@ -383,8 +364,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__select_chain_inf_nan_1(i1 %cond, define nofpclass(inf) float @ret_nofpclass_inf__fabs_select_ninf_rhs(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__fabs_select_ninf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0xFFF0000000000000 -; CHECK-NEXT: [[FABS:%.*]] = call float @llvm.fabs.f32(float [[SELECT]]) +; CHECK-NEXT: [[FABS:%.*]] = call float @llvm.fabs.f32(float [[X]]) ; CHECK-NEXT: ret float [[FABS]] ; %select = select i1 %cond, float %x, float 0xFFF0000000000000 @@ -396,8 +376,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__fabs_select_ninf_rhs(i1 %cond, f define nofpclass(inf) float @ret_nofpclass_inf__fabs_select_pinf_rhs(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__fabs_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[FABS:%.*]] = call float @llvm.fabs.f32(float [[SELECT]]) +; CHECK-NEXT: [[FABS:%.*]] = call float @llvm.fabs.f32(float [[X]]) ; CHECK-NEXT: ret float [[FABS]] ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 @@ -421,8 +400,7 @@ define nofpclass(ninf nnorm nsub nzero) float @ret_nofpclass_no_negatives__fabs_ define nofpclass(pinf pnorm psub pzero) float @ret_nofpclass_no_positives__fabs_select_pinf_rhs(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) float @ret_nofpclass_no_positives__fabs_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[FABS:%.*]] = call float @llvm.fabs.f32(float [[SELECT]]) +; CHECK-NEXT: [[FABS:%.*]] = call float @llvm.fabs.f32(float [[X]]) ; CHECK-NEXT: ret float [[FABS]] ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 @@ -446,9 +424,7 @@ define nofpclass(nan ninf nnorm nsub nzero) float @ret_nofpclass_no_negatives_na define nofpclass(nan pinf pnorm psub pzero) float @ret_nofpclass_no_positives_nan__fabs_select_pinf_rhs(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(nan pinf pzero psub pnorm) float @ret_nofpclass_no_positives_nan__fabs_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[FABS:%.*]] = call float @llvm.fabs.f32(float [[SELECT]]) -; CHECK-NEXT: ret float [[FABS]] +; CHECK-NEXT: ret float poison ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 %fabs = call float @llvm.fabs.f32(float %select) @@ -459,8 +435,7 @@ define nofpclass(nan pinf pnorm psub pzero) float @ret_nofpclass_no_positives_na define nofpclass(inf) float @ret_nofpclass_inf__fneg_select_ninf_rhs(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__fneg_select_ninf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0xFFF0000000000000 -; CHECK-NEXT: [[FNEG:%.*]] = fneg float [[SELECT]] +; CHECK-NEXT: [[FNEG:%.*]] = fneg float [[X]] ; CHECK-NEXT: ret float [[FNEG]] ; %select = select i1 %cond, float %x, float 0xFFF0000000000000 @@ -472,8 +447,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__fneg_select_ninf_rhs(i1 %cond, f define nofpclass(inf nnorm nsub nzero) float @ret_nofpclass_nonegatives_noinf___fneg_select_pinf_rhs(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(inf nzero nsub nnorm) float @ret_nofpclass_nonegatives_noinf___fneg_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[FNEG:%.*]] = fneg float [[SELECT]] +; CHECK-NEXT: [[FNEG:%.*]] = fneg float [[X]] ; CHECK-NEXT: ret float [[FNEG]] ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 @@ -485,8 +459,7 @@ define nofpclass(inf nnorm nsub nzero) float @ret_nofpclass_nonegatives_noinf___ define nofpclass(inf nnorm nsub nzero) float @ret_nofpclass_nonegatives_noinf___fneg_select_ninf_lhs(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(inf nzero nsub nnorm) float @ret_nofpclass_nonegatives_noinf___fneg_select_ninf_lhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float 0xFFF0000000000000, float [[X]] -; CHECK-NEXT: [[FNEG:%.*]] = fneg float [[SELECT]] +; CHECK-NEXT: [[FNEG:%.*]] = fneg float [[X]] ; CHECK-NEXT: ret float [[FNEG]] ; %select = select i1 %cond, float 0xFFF0000000000000, float %x @@ -510,8 +483,7 @@ define nofpclass(pzero psub pnorm pinf) float @ret_nofpclass_nopositives___fneg_ define nofpclass(inf) float @ret_nofpclass_inf__fneg_fabs_select_pinf_rhs(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__fneg_fabs_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[FABS:%.*]] = call float @llvm.fabs.f32(float [[SELECT]]) +; CHECK-NEXT: [[FABS:%.*]] = call float @llvm.fabs.f32(float [[X]]) ; CHECK-NEXT: [[FNEG:%.*]] = fneg float [[FABS]] ; CHECK-NEXT: ret float [[FNEG]] ; @@ -525,8 +497,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__fneg_fabs_select_pinf_rhs(i1 %co define nofpclass(ninf nnorm nsub nzero) float @ret_nofpclass_nonegatives__fneg_fabs_select_pinf_rhs(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) float @ret_nofpclass_nonegatives__fneg_fabs_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[FABS:%.*]] = call float @llvm.fabs.f32(float [[SELECT]]) +; CHECK-NEXT: [[FABS:%.*]] = call float @llvm.fabs.f32(float [[X]]) ; CHECK-NEXT: [[FNEG:%.*]] = fneg float [[FABS]] ; CHECK-NEXT: ret float [[FNEG]] ; @@ -541,10 +512,7 @@ define nofpclass(ninf nnorm nsub nzero) float @ret_nofpclass_nonegatives__fneg_f define nofpclass(nan ninf nnorm nsub nzero) float @ret_nofpclass_nonegatives_nonan__fneg_fabs_select_pinf_rhs(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(nan ninf nzero nsub nnorm) float @ret_nofpclass_nonegatives_nonan__fneg_fabs_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[FABS:%.*]] = call float @llvm.fabs.f32(float [[SELECT]]) -; CHECK-NEXT: [[FNEG:%.*]] = fneg float [[FABS]] -; CHECK-NEXT: ret float [[FNEG]] +; CHECK-NEXT: ret float poison ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 %fabs = call float @llvm.fabs.f32(float %select) @@ -556,8 +524,7 @@ define nofpclass(nan ninf nnorm nsub nzero) float @ret_nofpclass_nonegatives_non define nofpclass(inf) float @ret_nofpclass_inf__copysign_unknown_select_pinf_rhs(i1 %cond, float %x, float %unknown.sign) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__copysign_unknown_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]], float [[UNKNOWN_SIGN:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.copysign.f32(float [[SELECT]], float [[UNKNOWN_SIGN]]) +; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.copysign.f32(float [[X]], float [[UNKNOWN_SIGN]]) ; CHECK-NEXT: ret float [[COPYSIGN]] ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 @@ -568,8 +535,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__copysign_unknown_select_pinf_rhs define nofpclass(inf) float @ret_nofpclass_inf__copysign_positive_select_pinf_rhs(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__copysign_positive_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.fabs.f32(float [[SELECT]]) +; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.fabs.f32(float [[X]]) ; CHECK-NEXT: ret float [[COPYSIGN]] ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 @@ -580,8 +546,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__copysign_positive_select_pinf_rh define nofpclass(inf) float @ret_nofpclass_inf__copysign_negative_select_pinf_rhs(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__copysign_negative_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[TMP1:%.*]] = call float @llvm.fabs.f32(float [[SELECT]]) +; CHECK-NEXT: [[TMP1:%.*]] = call float @llvm.fabs.f32(float [[X]]) ; CHECK-NEXT: [[COPYSIGN:%.*]] = fneg float [[TMP1]] ; CHECK-NEXT: ret float [[COPYSIGN]] ; @@ -594,7 +559,8 @@ define nofpclass(inf) float @ret_nofpclass_inf__copysign_negative_select_pinf_rh define nofpclass(pinf pnorm psub pzero) float @ret_nofpclass_nopositives_copysign(float %x, float %unknown.sign) { ; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) float @ret_nofpclass_nopositives_copysign ; CHECK-SAME: (float [[X:%.*]], float [[UNKNOWN_SIGN:%.*]]) { -; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.copysign.f32(float [[X]], float [[UNKNOWN_SIGN]]) +; CHECK-NEXT: [[TMP1:%.*]] = call float @llvm.fabs.f32(float [[X]]) +; CHECK-NEXT: [[COPYSIGN:%.*]] = fneg float [[TMP1]] ; CHECK-NEXT: ret float [[COPYSIGN]] ; %copysign = call float @llvm.copysign.f32(float %x, float %unknown.sign) @@ -605,7 +571,8 @@ define nofpclass(pinf pnorm psub pzero) float @ret_nofpclass_nopositives_copysig define nofpclass(pinf pnorm psub pzero) float @ret_nofpclass_nopositives_copysign_nnan_flag(float %x, float %unknown.sign) { ; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) float @ret_nofpclass_nopositives_copysign_nnan_flag ; CHECK-SAME: (float [[X:%.*]], float [[UNKNOWN_SIGN:%.*]]) { -; CHECK-NEXT: [[COPYSIGN:%.*]] = call nnan float @llvm.copysign.f32(float [[X]], float [[UNKNOWN_SIGN]]) +; CHECK-NEXT: [[TMP1:%.*]] = call nnan float @llvm.fabs.f32(float [[X]]) +; CHECK-NEXT: [[COPYSIGN:%.*]] = fneg nnan float [[TMP1]] ; CHECK-NEXT: ret float [[COPYSIGN]] ; %copysign = call nnan float @llvm.copysign.f32(float %x, float %unknown.sign) @@ -616,7 +583,8 @@ define nofpclass(pinf pnorm psub pzero) float @ret_nofpclass_nopositives_copysig define nofpclass(nan pinf pnorm psub pzero) float @ret_nofpclass_nopositives_nonan_copysign(float %x, float %unknown.sign) { ; CHECK-LABEL: define nofpclass(nan pinf pzero psub pnorm) float @ret_nofpclass_nopositives_nonan_copysign ; CHECK-SAME: (float [[X:%.*]], float [[UNKNOWN_SIGN:%.*]]) { -; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.copysign.f32(float [[X]], float [[UNKNOWN_SIGN]]) +; CHECK-NEXT: [[TMP1:%.*]] = call float @llvm.fabs.f32(float [[X]]) +; CHECK-NEXT: [[COPYSIGN:%.*]] = fneg float [[TMP1]] ; CHECK-NEXT: ret float [[COPYSIGN]] ; %copysign = call float @llvm.copysign.f32(float %x, float %unknown.sign) @@ -627,7 +595,7 @@ define nofpclass(nan pinf pnorm psub pzero) float @ret_nofpclass_nopositives_non define nofpclass(ninf nnorm nsub nzero) float @ret_nofpclass_nonegatives_copysign(float %x, float %unknown.sign) { ; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) float @ret_nofpclass_nonegatives_copysign ; CHECK-SAME: (float [[X:%.*]], float [[UNKNOWN_SIGN:%.*]]) { -; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.copysign.f32(float [[X]], float [[UNKNOWN_SIGN]]) +; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.fabs.f32(float [[X]]) ; CHECK-NEXT: ret float [[COPYSIGN]] ; %copysign = call float @llvm.copysign.f32(float %x, float %unknown.sign) @@ -638,7 +606,7 @@ define nofpclass(ninf nnorm nsub nzero) float @ret_nofpclass_nonegatives_copysig define nofpclass(ninf nnorm nsub nzero) float @ret_nofpclass_nonegatives_copysign_nnan_flag(float %x, float %unknown.sign) { ; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) float @ret_nofpclass_nonegatives_copysign_nnan_flag ; CHECK-SAME: (float [[X:%.*]], float [[UNKNOWN_SIGN:%.*]]) { -; CHECK-NEXT: [[COPYSIGN:%.*]] = call nnan float @llvm.copysign.f32(float [[X]], float [[UNKNOWN_SIGN]]) +; CHECK-NEXT: [[COPYSIGN:%.*]] = call nnan float @llvm.fabs.f32(float [[X]]) ; CHECK-NEXT: ret float [[COPYSIGN]] ; %copysign = call nnan float @llvm.copysign.f32(float %x, float %unknown.sign) @@ -649,7 +617,7 @@ define nofpclass(ninf nnorm nsub nzero) float @ret_nofpclass_nonegatives_copysig define nofpclass(nan ninf nnorm nsub nzero) float @ret_nofpclass_nonegatives_nonan_copysign(float %x, float %unknown.sign) { ; CHECK-LABEL: define nofpclass(nan ninf nzero nsub nnorm) float @ret_nofpclass_nonegatives_nonan_copysign ; CHECK-SAME: (float [[X:%.*]], float [[UNKNOWN_SIGN:%.*]]) { -; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.copysign.f32(float [[X]], float [[UNKNOWN_SIGN]]) +; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.fabs.f32(float [[X]]) ; CHECK-NEXT: ret float [[COPYSIGN]] ; %copysign = call float @llvm.copysign.f32(float %x, float %unknown.sign) @@ -659,8 +627,7 @@ define nofpclass(nan ninf nnorm nsub nzero) float @ret_nofpclass_nonegatives_non define nofpclass(pinf pnorm psub pzero) float @ret_nofpclass_nopositives__copysign_fabs_select_pinf_rhs(i1 %cond, float %x, float %sign) { ; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) float @ret_nofpclass_nopositives__copysign_fabs_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]], float [[SIGN:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.fabs.f32(float [[SELECT]]) +; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.fabs.f32(float [[X]]) ; CHECK-NEXT: ret float [[COPYSIGN]] ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 @@ -673,8 +640,7 @@ define nofpclass(pinf pnorm psub pzero) float @ret_nofpclass_nopositives__copysi define nofpclass(inf nnorm nsub nzero) float @ret_nofpclass_no_negatives_noinf__copysign_unknown_select_pinf_rhs(i1 %cond, float %x, float %unknown.sign) { ; CHECK-LABEL: define nofpclass(inf nzero nsub nnorm) float @ret_nofpclass_no_negatives_noinf__copysign_unknown_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]], float [[UNKNOWN_SIGN:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.copysign.f32(float [[SELECT]], float [[UNKNOWN_SIGN]]) +; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.fabs.f32(float [[X]]) ; CHECK-NEXT: ret float [[COPYSIGN]] ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 @@ -686,8 +652,8 @@ define nofpclass(inf nnorm nsub nzero) float @ret_nofpclass_no_negatives_noinf__ define nofpclass(inf pnorm psub pzero) float @ret_nofpclass_no_positives_noinf__copysign_unknown_select_pinf_rhs(i1 %cond, float %x, float %unknown.sign) { ; CHECK-LABEL: define nofpclass(inf pzero psub pnorm) float @ret_nofpclass_no_positives_noinf__copysign_unknown_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]], float [[UNKNOWN_SIGN:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.copysign.f32(float [[SELECT]], float [[UNKNOWN_SIGN]]) +; CHECK-NEXT: [[TMP1:%.*]] = call float @llvm.fabs.f32(float [[X]]) +; CHECK-NEXT: [[COPYSIGN:%.*]] = fneg float [[TMP1]] ; CHECK-NEXT: ret float [[COPYSIGN]] ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 @@ -700,7 +666,7 @@ define nofpclass(ninf nnorm nsub nzero) float @ret_nofpclass_no_negatives__copys ; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) float @ret_nofpclass_no_negatives__copysign_unknown_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]], float [[UNKNOWN_SIGN:%.*]]) { ; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.copysign.f32(float [[SELECT]], float [[UNKNOWN_SIGN]]) +; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.fabs.f32(float [[SELECT]]) ; CHECK-NEXT: ret float [[COPYSIGN]] ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 @@ -713,7 +679,8 @@ define nofpclass(pinf pnorm psub pzero) float @ret_nofpclass_no_positives__copys ; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) float @ret_nofpclass_no_positives__copysign_unknown_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]], float [[UNKNOWN_SIGN:%.*]]) { ; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.copysign.f32(float [[SELECT]], float [[UNKNOWN_SIGN]]) +; CHECK-NEXT: [[TMP1:%.*]] = call float @llvm.fabs.f32(float [[SELECT]]) +; CHECK-NEXT: [[COPYSIGN:%.*]] = fneg float [[TMP1]] ; CHECK-NEXT: ret float [[COPYSIGN]] ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 @@ -726,7 +693,7 @@ define nofpclass(nan ninf nnorm nsub nzero) float @ret_nofpclass_no_negatives_no ; CHECK-LABEL: define nofpclass(nan ninf nzero nsub nnorm) float @ret_nofpclass_no_negatives_nonan__copysign_unknown_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]], float [[UNKNOWN_SIGN:%.*]]) { ; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.copysign.f32(float [[SELECT]], float [[UNKNOWN_SIGN]]) +; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.fabs.f32(float [[SELECT]]) ; CHECK-NEXT: ret float [[COPYSIGN]] ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 @@ -739,7 +706,8 @@ define nofpclass(nan pinf pnorm psub pzero) float @ret_nofpclass_no_positives_no ; CHECK-LABEL: define nofpclass(nan pinf pzero psub pnorm) float @ret_nofpclass_no_positives_nonan__copysign_unknown_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]], float [[UNKNOWN_SIGN:%.*]]) { ; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[COPYSIGN:%.*]] = call float @llvm.copysign.f32(float [[SELECT]], float [[UNKNOWN_SIGN]]) +; CHECK-NEXT: [[TMP1:%.*]] = call float @llvm.fabs.f32(float [[SELECT]]) +; CHECK-NEXT: [[COPYSIGN:%.*]] = fneg float [[TMP1]] ; CHECK-NEXT: ret float [[COPYSIGN]] ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 @@ -790,9 +758,7 @@ define nofpclass(nan ninf nnorm nsub nzero) float @ret_nofpclass_nan_negatives__ define nofpclass(nan ninf nnorm nsub zero) float @ret_nofpclass_nan_negatives_zero__select_clamp_pos_to_zero(float %x) { ; CHECK-LABEL: define nofpclass(nan ninf zero nsub nnorm) float @ret_nofpclass_nan_negatives_zero__select_clamp_pos_to_zero ; CHECK-SAME: (float [[X:%.*]]) { -; CHECK-NEXT: [[IS_GT_ZERO:%.*]] = fcmp ogt float [[X]], 0.000000e+00 -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[IS_GT_ZERO]], float 0.000000e+00, float [[X]] -; CHECK-NEXT: ret float [[SELECT]] +; CHECK-NEXT: ret float [[X]] ; %is.gt.zero = fcmp ogt float %x, 0.0 %select = select i1 %is.gt.zero, float 0.0, float %x @@ -803,9 +769,7 @@ define nofpclass(nan ninf nnorm nsub zero) float @ret_nofpclass_nan_negatives_ze define nofpclass(ninf nnorm nsub zero) float @ret_nofpclass_negatives_zero__select_clamp_pos_to_zero(float %x) { ; CHECK-LABEL: define nofpclass(ninf zero nsub nnorm) float @ret_nofpclass_negatives_zero__select_clamp_pos_to_zero ; CHECK-SAME: (float [[X:%.*]]) { -; CHECK-NEXT: [[IS_GT_ZERO:%.*]] = fcmp ogt float [[X]], 0.000000e+00 -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[IS_GT_ZERO]], float 0.000000e+00, float [[X]] -; CHECK-NEXT: ret float [[SELECT]] +; CHECK-NEXT: ret float [[X]] ; %is.gt.zero = fcmp ogt float %x, 0.0 %select = select i1 %is.gt.zero, float 0.0, float %x @@ -819,8 +783,7 @@ define nofpclass(inf) float @ret_nofpclass_noinfs__assumed_isinf__select_pinf_lh ; CHECK-NEXT: [[FABS_X:%.*]] = call float @llvm.fabs.f32(float [[X]]) ; CHECK-NEXT: [[X_IS_INF:%.*]] = fcmp oeq float [[FABS_X]], 0x7FF0000000000000 ; CHECK-NEXT: call void @llvm.assume(i1 [[X_IS_INF]]) -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float [[Y]] -; CHECK-NEXT: ret float [[SELECT]] +; CHECK-NEXT: ret float [[Y]] ; %fabs.x = call float @llvm.fabs.f32(float %x) %x.is.inf = fcmp oeq float %fabs.x, 0x7FF0000000000000 @@ -838,18 +801,13 @@ define nofpclass(nan inf nzero nsub nnorm) float @powr_issue64870(float nofpclas ; CHECK-NEXT: [[I1:%.*]] = tail call float @llvm.log2.f32(float [[I]]) ; CHECK-NEXT: [[I2:%.*]] = fmul float [[I1]], [[Y]] ; CHECK-NEXT: [[I3:%.*]] = tail call nofpclass(ninf nzero nsub nnorm) float @llvm.exp2.f32(float [[I2]]) -; CHECK-NEXT: [[I4:%.*]] = fcmp olt float [[Y]], 0.000000e+00 -; CHECK-NEXT: [[I5:%.*]] = select i1 [[I4]], float 0x7FF0000000000000, float 0.000000e+00 ; CHECK-NEXT: [[I6:%.*]] = fcmp oeq float [[X]], 0.000000e+00 -; CHECK-NEXT: [[I7:%.*]] = select i1 [[I6]], float [[I5]], float [[I3]] +; CHECK-NEXT: [[I7:%.*]] = select i1 [[I6]], float 0.000000e+00, float [[I3]] ; CHECK-NEXT: [[I8:%.*]] = fcmp oeq float [[Y]], 0.000000e+00 -; CHECK-NEXT: [[I9:%.*]] = select i1 [[I6]], float 0x7FF8000000000000, float 1.000000e+00 -; CHECK-NEXT: [[I10:%.*]] = select i1 [[I8]], float [[I9]], float [[I7]] ; CHECK-NEXT: [[I11:%.*]] = fcmp oeq float [[X]], 1.000000e+00 -; CHECK-NEXT: [[I12:%.*]] = select i1 [[I11]], float 1.000000e+00, float [[I10]] -; CHECK-NEXT: [[I13:%.*]] = fcmp olt float [[X]], 0.000000e+00 -; CHECK-NEXT: [[I14:%.*]] = select i1 [[I13]], float 0x7FF8000000000000, float [[I12]] -; CHECK-NEXT: ret float [[I14]] +; CHECK-NEXT: [[TMP0:%.*]] = select i1 [[I11]], i1 true, i1 [[I8]] +; CHECK-NEXT: [[I12:%.*]] = select i1 [[TMP0]], float 1.000000e+00, float [[I7]] +; CHECK-NEXT: ret float [[I12]] ; entry: %i = tail call float @llvm.fabs.f32(float %x) @@ -881,12 +839,8 @@ define nofpclass(nan inf nzero nsub nnorm) float @test_powr_issue64870_2(float n ; CHECK-NEXT: [[I4:%.*]] = select i1 [[I]], float 0x7FF8000000000000, float [[ARG1]] ; CHECK-NEXT: [[I5:%.*]] = fmul float [[I4]], [[I3]] ; CHECK-NEXT: [[I6:%.*]] = tail call noundef nofpclass(ninf nzero nsub nnorm) float @llvm.exp2.f32(float noundef [[I5]]) -; CHECK-NEXT: [[I7:%.*]] = fcmp olt float [[I4]], 0.000000e+00 -; CHECK-NEXT: [[I8:%.*]] = select i1 [[I7]], float 0x7FF0000000000000, float 0.000000e+00 -; CHECK-NEXT: [[I9:%.*]] = fcmp ueq float [[I4]], 0.000000e+00 ; CHECK-NEXT: [[I10:%.*]] = fcmp oeq float [[I2]], 0.000000e+00 -; CHECK-NEXT: [[I11:%.*]] = select i1 [[I9]], float 0x7FF8000000000000, float [[I8]] -; CHECK-NEXT: [[I12:%.*]] = select i1 [[I10]], float [[I11]], float [[I6]] +; CHECK-NEXT: [[I12:%.*]] = select i1 [[I10]], float 0.000000e+00, float [[I6]] ; CHECK-NEXT: ret float [[I12]] ; bb: @@ -923,16 +877,10 @@ define nofpclass(nan inf) float @pow_f32(float nofpclass(nan inf) %arg, float no ; CHECK-NEXT: [[I11:%.*]] = and i1 [[I7]], [[I10]] ; CHECK-NEXT: [[I12:%.*]] = select i1 [[I11]], float [[ARG]], float 1.000000e+00 ; CHECK-NEXT: [[I13:%.*]] = tail call noundef float @llvm.copysign.f32(float noundef [[I4]], float noundef [[I12]]) -; CHECK-NEXT: [[I14:%.*]] = fcmp olt float [[ARG]], 0.000000e+00 -; CHECK-NEXT: [[I15:%.*]] = select i1 [[I7]], float [[I13]], float 0x7FF8000000000000 -; CHECK-NEXT: [[I16:%.*]] = select i1 [[I14]], float [[I15]], float [[I13]] ; CHECK-NEXT: [[I17:%.*]] = fcmp oeq float [[ARG]], 0.000000e+00 -; CHECK-NEXT: [[I18:%.*]] = fcmp olt float [[ARG1]], 0.000000e+00 -; CHECK-NEXT: [[I19:%.*]] = xor i1 [[I17]], [[I18]] -; CHECK-NEXT: [[I20:%.*]] = select i1 [[I19]], float 0.000000e+00, float 0x7FF0000000000000 ; CHECK-NEXT: [[I21:%.*]] = select i1 [[I11]], float [[ARG]], float 0.000000e+00 -; CHECK-NEXT: [[I22:%.*]] = tail call noundef nofpclass(nan sub norm) float @llvm.copysign.f32(float noundef [[I20]], float noundef [[I21]]) -; CHECK-NEXT: [[I23:%.*]] = select i1 [[I17]], float [[I22]], float [[I16]] +; CHECK-NEXT: [[I22:%.*]] = tail call noundef nofpclass(nan sub norm) float @llvm.copysign.f32(float noundef 0.000000e+00, float noundef [[I21]]) +; CHECK-NEXT: [[I23:%.*]] = select i1 [[I17]], float [[I22]], float [[I13]] ; CHECK-NEXT: [[I24:%.*]] = fcmp oeq float [[ARG]], 1.000000e+00 ; CHECK-NEXT: [[I25:%.*]] = fcmp oeq float [[ARG1]], 0.000000e+00 ; CHECK-NEXT: [[I26:%.*]] = or i1 [[I24]], [[I25]] @@ -977,8 +925,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__select_nofpclass_call_only_inf(i ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__select_nofpclass_call_only_inf ; CHECK-SAME: (i1 [[COND:%.*]], float [[Y:%.*]]) { ; CHECK-NEXT: [[MUST_BE_INF:%.*]] = call nofpclass(nan zero sub norm) float @extern() -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[MUST_BE_INF]], float [[Y]] -; CHECK-NEXT: ret float [[SELECT]] +; CHECK-NEXT: ret float [[Y]] ; %must.be.inf = call nofpclass(nan norm zero sub) float @extern() %select = select i1 %cond, float %must.be.inf, float %y @@ -989,7 +936,7 @@ define nofpclass(pinf) float @ret_nofpclass_pinf__nofpclass_call_only_inf(i1 %co ; CHECK-LABEL: define nofpclass(pinf) float @ret_nofpclass_pinf__nofpclass_call_only_inf ; CHECK-SAME: (i1 [[COND:%.*]], float [[Y:%.*]]) { ; CHECK-NEXT: [[MUST_BE_INF:%.*]] = call nofpclass(nan zero sub norm) float @extern() -; CHECK-NEXT: ret float [[MUST_BE_INF]] +; CHECK-NEXT: ret float 0xFFF0000000000000 ; %must.be.inf = call nofpclass(nan norm zero sub) float @extern() ret float %must.be.inf @@ -999,7 +946,7 @@ define nofpclass(ninf) float @ret_nofpclass_ninf__nofpclass_call_only_inf(i1 %co ; CHECK-LABEL: define nofpclass(ninf) float @ret_nofpclass_ninf__nofpclass_call_only_inf ; CHECK-SAME: (i1 [[COND:%.*]], float [[Y:%.*]]) { ; CHECK-NEXT: [[MUST_BE_INF:%.*]] = call nofpclass(nan zero sub norm) float @extern() -; CHECK-NEXT: ret float [[MUST_BE_INF]] +; CHECK-NEXT: ret float 0x7FF0000000000000 ; %must.be.inf = call nofpclass(nan norm zero sub) float @extern() ret float %must.be.inf @@ -1009,7 +956,7 @@ define nofpclass(nzero) float @ret_nofpclass_nzero__nofpclass_call_only_zero(i1 ; CHECK-LABEL: define nofpclass(nzero) float @ret_nofpclass_nzero__nofpclass_call_only_zero ; CHECK-SAME: (i1 [[COND:%.*]], float [[Y:%.*]]) { ; CHECK-NEXT: [[MUST_BE_ZERO:%.*]] = call nofpclass(nan inf sub norm) float @extern() -; CHECK-NEXT: ret float [[MUST_BE_ZERO]] +; CHECK-NEXT: ret float 0.000000e+00 ; %must.be.zero = call nofpclass(nan sub norm inf) float @extern() ret float %must.be.zero @@ -1019,7 +966,7 @@ define nofpclass(pzero) float @ret_nofpclass_pzero__nofpclass_call_only_zero(i1 ; CHECK-LABEL: define nofpclass(pzero) float @ret_nofpclass_pzero__nofpclass_call_only_zero ; CHECK-SAME: (i1 [[COND:%.*]], float [[Y:%.*]]) { ; CHECK-NEXT: [[MUST_BE_ZERO:%.*]] = call nofpclass(nan inf sub norm) float @extern() -; CHECK-NEXT: ret float [[MUST_BE_ZERO]] +; CHECK-NEXT: ret float -0.000000e+00 ; %must.be.zero = call nofpclass(nan sub norm inf) float @extern() ret float %must.be.zero @@ -1133,8 +1080,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__recursive_phi_0(i1 %cond0, float ; CHECK-NEXT: [[LOOP_COND:%.*]] = call i1 @loop.cond() ; CHECK-NEXT: br i1 [[LOOP_COND]], label [[RET]], label [[LOOP]] ; CHECK: ret: -; CHECK-NEXT: [[PHI_RET:%.*]] = phi float [ 0.000000e+00, [[ENTRY:%.*]] ], [ 0x7FF0000000000000, [[LOOP]] ] -; CHECK-NEXT: ret float [[PHI_RET]] +; CHECK-NEXT: ret float 0.000000e+00 ; entry: br i1 %cond0, label %loop, label %ret @@ -1159,7 +1105,7 @@ define nofpclass(inf) float @ret_nofpclass_inf__recursive_phi_1(i1 %cond0, float ; CHECK-NEXT: [[LOOP_COND:%.*]] = call i1 @loop.cond() ; CHECK-NEXT: br i1 [[LOOP_COND]], label [[RET]], label [[LOOP]] ; CHECK: ret: -; CHECK-NEXT: ret float 0x7FF0000000000000 +; CHECK-NEXT: ret float poison ; entry: br i1 %cond0, label %loop, label %ret @@ -1211,8 +1157,7 @@ ret: define nofpclass(inf) float @ret_nofpclass_inf__arithmetic_fence_select_pinf_rhs(i1 %cond, float %x) { ; CHECK-LABEL: define nofpclass(inf) float @ret_nofpclass_inf__arithmetic_fence_select_pinf_rhs ; CHECK-SAME: (i1 [[COND:%.*]], float [[X:%.*]]) { -; CHECK-NEXT: [[SELECT:%.*]] = select i1 [[COND]], float [[X]], float 0x7FF0000000000000 -; CHECK-NEXT: [[FENCE:%.*]] = call float @llvm.arithmetic.fence.f32(float [[SELECT]]) +; CHECK-NEXT: [[FENCE:%.*]] = call float @llvm.arithmetic.fence.f32(float [[X]]) ; CHECK-NEXT: ret float [[FENCE]] ; %select = select i1 %cond, float %x, float 0x7FF0000000000000 From 6407fe78971dacfecefeaacfa99d24b94601463c Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Mon, 29 Jan 2024 14:20:42 +0530 Subject: [PATCH 2/4] Enable only under off-by-default flag for now --- llvm/lib/Transforms/InstCombine/InstructionCombining.cpp | 9 +++++++++ .../Transforms/InstCombine/simplify-demanded-fpclass.ll | 2 +- 2 files changed, 10 insertions(+), 1 deletion(-) diff --git a/llvm/lib/Transforms/InstCombine/InstructionCombining.cpp b/llvm/lib/Transforms/InstCombine/InstructionCombining.cpp index a45f9595a9907..b9f7e38731b51 100644 --- a/llvm/lib/Transforms/InstCombine/InstructionCombining.cpp +++ b/llvm/lib/Transforms/InstCombine/InstructionCombining.cpp @@ -142,6 +142,12 @@ static cl::opt MaxArraySize("instcombine-maxarray-size", cl::init(1024), cl::desc("Maximum array size considered when doing a combine")); +// TODO: Remove this option +static cl::opt EnableSimplifyDemandedUseFPClass( + "instcombine-simplify-demanded-fp-class", + cl::desc("Enable demanded floating-point class optimizations"), + cl::init(false)); + // FIXME: Remove this flag when it is no longer necessary to convert // llvm.dbg.declare to avoid inaccurate debug info. Setting this to false // increases variable availability at the cost of accuracy. Variables that @@ -3093,6 +3099,9 @@ Instruction *InstCombinerImpl::visitFree(CallInst &FI, Value *Op) { } Instruction *InstCombinerImpl::visitReturnInst(ReturnInst &RI) { + if (!EnableSimplifyDemandedUseFPClass) + return nullptr; + Value *RetVal = RI.getReturnValue(); if (!RetVal || !AttributeFuncs::isNoFPClassCompatibleType(RetVal->getType())) return nullptr; diff --git a/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass.ll b/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass.ll index 4f9396add2370..9a90c97da90a1 100644 --- a/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass.ll +++ b/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass.ll @@ -1,5 +1,5 @@ ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 2 -; RUN: opt -S -passes=instcombine < %s | FileCheck %s +; RUN: opt -S -passes=instcombine -instcombine-simplify-demanded-fp-class < %s | FileCheck %s declare float @llvm.fabs.f32(float) declare float @llvm.copysign.f32(float, float) From d2fa818e4a61646e7073be90a94a599503753668 Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Thu, 8 Feb 2024 11:38:02 +0530 Subject: [PATCH 3/4] update_test_checks formatting change --- llvm/test/Transforms/InstCombine/simplify-demanded-fpclass.ll | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass.ll b/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass.ll index 9a90c97da90a1..dd9b71415bd6d 100644 --- a/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass.ll +++ b/llvm/test/Transforms/InstCombine/simplify-demanded-fpclass.ll @@ -1126,8 +1126,8 @@ define nofpclass(inf) float @ret_nofpclass_inf__phi_switch_repeated_predecessor( ; CHECK-SAME: (i32 [[SWITCH:%.*]], float [[UNKNOWN:%.*]]) { ; CHECK-NEXT: entry: ; CHECK-NEXT: switch i32 [[SWITCH]], label [[RET:%.*]] [ -; CHECK-NEXT: i32 0, label [[LOOP:%.*]] -; CHECK-NEXT: i32 1, label [[LOOP]] +; CHECK-NEXT: i32 0, label [[LOOP:%.*]] +; CHECK-NEXT: i32 1, label [[LOOP]] ; CHECK-NEXT: ] ; CHECK: loop: ; CHECK-NEXT: [[PHI_LOOP:%.*]] = phi float [ 0x7FF0000000000000, [[ENTRY:%.*]] ], [ 0x7FF0000000000000, [[ENTRY]] ], [ [[UNKNOWN]], [[LOOP]] ] From 04a5371c3f806209225a2d1bf0e8f6b304934daf Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Thu, 8 Feb 2024 13:29:57 +0530 Subject: [PATCH 4/4] Update clang test since it's now off by default --- clang/test/Headers/__clang_hip_math.hip | 56 ++++++------------------- 1 file changed, 12 insertions(+), 44 deletions(-) diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index 37099de74fb8e..e9a9cb4bb3c74 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -2557,65 +2557,33 @@ extern "C" __device__ double test_nan(const char *tag) { return nan(tag); } -// DEFAULT-LABEL: @test_nanf_emptystr( -// DEFAULT-NEXT: entry: -// DEFAULT-NEXT: ret float 0x7FF8000000000000 -// -// FINITEONLY-LABEL: @test_nanf_emptystr( -// FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: ret float poison -// -// APPROX-LABEL: @test_nanf_emptystr( -// APPROX-NEXT: entry: -// APPROX-NEXT: ret float 0x7FF8000000000000 +// CHECK-LABEL: @test_nanf_emptystr( +// CHECK-NEXT: entry: +// CHECK-NEXT: ret float 0x7FF8000000000000 // extern "C" __device__ float test_nanf_emptystr() { return nanf(""); } -// DEFAULT-LABEL: @test_nan_emptystr( -// DEFAULT-NEXT: entry: -// DEFAULT-NEXT: ret double 0x7FF8000000000000 -// -// FINITEONLY-LABEL: @test_nan_emptystr( -// FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: ret double poison -// -// APPROX-LABEL: @test_nan_emptystr( -// APPROX-NEXT: entry: -// APPROX-NEXT: ret double 0x7FF8000000000000 +// CHECK-LABEL: @test_nan_emptystr( +// CHECK-NEXT: entry: +// CHECK-NEXT: ret double 0x7FF8000000000000 // extern "C" __device__ double test_nan_emptystr() { return nan(""); } -// DEFAULT-LABEL: @test_nanf_fill( -// DEFAULT-NEXT: entry: -// DEFAULT-NEXT: ret float 0x7FF8000000000000 -// -// FINITEONLY-LABEL: @test_nanf_fill( -// FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: ret float poison -// -// APPROX-LABEL: @test_nanf_fill( -// APPROX-NEXT: entry: -// APPROX-NEXT: ret float 0x7FF8000000000000 +// CHECK-LABEL: @test_nanf_fill( +// CHECK-NEXT: entry: +// CHECK-NEXT: ret float 0x7FF8000000000000 // extern "C" __device__ float test_nanf_fill() { return nanf("0x456"); } -// DEFAULT-LABEL: @test_nan_fill( -// DEFAULT-NEXT: entry: -// DEFAULT-NEXT: ret double 0x7FF8000000000000 -// -// FINITEONLY-LABEL: @test_nan_fill( -// FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: ret double poison -// -// APPROX-LABEL: @test_nan_fill( -// APPROX-NEXT: entry: -// APPROX-NEXT: ret double 0x7FF8000000000000 +// CHECK-LABEL: @test_nan_fill( +// CHECK-NEXT: entry: +// CHECK-NEXT: ret double 0x7FF8000000000000 // extern "C" __device__ double test_nan_fill() { return nan("0x123");