https://github.com/AlexMaclean updated
https://github.com/llvm/llvm-project/pull/97762
>From c2913d1074c5bfa771379d68e9ba728a3d1d1ce5 Mon Sep 17 00:00:00 2001
From: Alex MacLean <amacl...@nvidia.com>
Date: Mon, 1 Jul 2024 17:06:56 +0000
Subject: [PATCH 1/4] [ValueTracking] use KnownBits to compute fpclass from
bitcast
---
llvm/lib/Analysis/ValueTracking.cpp | 30 ++++++
llvm/test/Transforms/Attributor/nofpclass.ll | 104 +++++++++++++++++++
2 files changed, 134 insertions(+)
diff --git a/llvm/lib/Analysis/ValueTracking.cpp
b/llvm/lib/Analysis/ValueTracking.cpp
index 85abf00774a02..a16c8e3d48403 100644
--- a/llvm/lib/Analysis/ValueTracking.cpp
+++ b/llvm/lib/Analysis/ValueTracking.cpp
@@ -5805,6 +5805,36 @@ void computeKnownFPClass(const Value *V, const APInt
&DemandedElts,
break;
}
+ case Instruction::BitCast: {
+ const Type *Ty = Op->getType();
+ const Value *Casted = Op->getOperand(0);
+ if (Ty->isVectorTy() || !Casted->getType()->isIntOrIntVectorTy())
+ break;
+
+ KnownBits Bits(Ty->getScalarSizeInBits());
+ computeKnownBits(Casted, Bits, Depth + 1, Q);
+
+ // Transfer information from the sign bit.
+ if (Bits.Zero.isSignBitSet())
+ Known.signBitMustBeZero();
+ else if (Bits.One.isSignBitSet())
+ Known.signBitMustBeOne();
+
+ if (Ty->isIEEE()) {
+ // IEEE floats are NaN when all bits of the exponent plus at least one of
+ // the fraction bits are 1. This means:
+ // - If we assume unknown bits are 0 and the value is NaN, it will
+ // always be NaN
+ // - If we assume unknown bits are 1 and the value is not NaN, it can
+ // never be NaN
+ if (APFloat(Ty->getFltSemantics(), Bits.One).isNaN())
+ Known.KnownFPClasses = fcNan;
+ else if (!APFloat(Ty->getFltSemantics(), ~Bits.Zero).isNaN())
+ Known.knownNot(fcNan);
+ }
+
+ break;
+ }
default:
break;
}
diff --git a/llvm/test/Transforms/Attributor/nofpclass.ll
b/llvm/test/Transforms/Attributor/nofpclass.ll
index 781ba636c3ab3..c5d562a436b33 100644
--- a/llvm/test/Transforms/Attributor/nofpclass.ll
+++ b/llvm/test/Transforms/Attributor/nofpclass.ll
@@ -2690,6 +2690,110 @@ entry:
ret double %abs
}
+define float @bitcast_to_float_sign_0(i32 %arg) {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
+; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) float
@bitcast_to_float_sign_0
+; CHECK-SAME: (i32 [[ARG:%.*]]) #[[ATTR3]] {
+; CHECK-NEXT: [[TMP1:%.*]] = lshr i32 [[ARG]], 1
+; CHECK-NEXT: [[TMP2:%.*]] = bitcast i32 [[TMP1]] to float
+; CHECK-NEXT: ret float [[TMP2]]
+;
+ %1 = lshr i32 %arg, 1
+ %2 = bitcast i32 %1 to float
+ ret float %2
+}
+
+define float @bitcast_to_float_nnan(i32 %arg) {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
+; CHECK-LABEL: define nofpclass(nan ninf nzero nsub nnorm) float
@bitcast_to_float_nnan
+; CHECK-SAME: (i32 [[ARG:%.*]]) #[[ATTR3]] {
+; CHECK-NEXT: [[TMP1:%.*]] = lshr i32 [[ARG]], 2
+; CHECK-NEXT: [[TMP2:%.*]] = bitcast i32 [[TMP1]] to float
+; CHECK-NEXT: ret float [[TMP2]]
+;
+ %1 = lshr i32 %arg, 2
+ %2 = bitcast i32 %1 to float
+ ret float %2
+}
+
+define float @bitcast_to_float_sign_1(i32 %arg) {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
+; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) float
@bitcast_to_float_sign_1
+; CHECK-SAME: (i32 [[ARG:%.*]]) #[[ATTR3]] {
+; CHECK-NEXT: [[TMP1:%.*]] = or i32 [[ARG]], -2147483648
+; CHECK-NEXT: [[TMP2:%.*]] = bitcast i32 [[TMP1]] to float
+; CHECK-NEXT: ret float [[TMP2]]
+;
+ %1 = or i32 %arg, -2147483648
+ %2 = bitcast i32 %1 to float
+ ret float %2
+}
+
+define float @bitcast_to_float_nan(i32 %arg) {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
+; CHECK-LABEL: define nofpclass(inf zero sub norm) float @bitcast_to_float_nan
+; CHECK-SAME: (i32 [[ARG:%.*]]) #[[ATTR3]] {
+; CHECK-NEXT: [[TMP1:%.*]] = or i32 [[ARG]], 2139095041
+; CHECK-NEXT: [[TMP2:%.*]] = bitcast i32 [[TMP1]] to float
+; CHECK-NEXT: ret float [[TMP2]]
+;
+ %1 = or i32 %arg, 2139095041
+ %2 = bitcast i32 %1 to float
+ ret float %2
+}
+
+define double @bitcast_to_double_sign_0(i64 %arg) {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
+; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) double
@bitcast_to_double_sign_0
+; CHECK-SAME: (i64 [[ARG:%.*]]) #[[ATTR3]] {
+; CHECK-NEXT: [[TMP1:%.*]] = lshr i64 [[ARG]], 1
+; CHECK-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to double
+; CHECK-NEXT: ret double [[TMP2]]
+;
+ %1 = lshr i64 %arg, 1
+ %2 = bitcast i64 %1 to double
+ ret double %2
+}
+
+define double @bitcast_to_double_nnan(i64 %arg) {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
+; CHECK-LABEL: define nofpclass(nan ninf nzero nsub nnorm) double
@bitcast_to_double_nnan
+; CHECK-SAME: (i64 [[ARG:%.*]]) #[[ATTR3]] {
+; CHECK-NEXT: [[TMP1:%.*]] = lshr i64 [[ARG]], 2
+; CHECK-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to double
+; CHECK-NEXT: ret double [[TMP2]]
+;
+ %1 = lshr i64 %arg, 2
+ %2 = bitcast i64 %1 to double
+ ret double %2
+}
+
+define double @bitcast_to_double_sign_1(i64 %arg) {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
+; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) double
@bitcast_to_double_sign_1
+; CHECK-SAME: (i64 [[ARG:%.*]]) #[[ATTR3]] {
+; CHECK-NEXT: [[TMP1:%.*]] = or i64 [[ARG]], -9223372036854775808
+; CHECK-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to double
+; CHECK-NEXT: ret double [[TMP2]]
+;
+ %1 = or i64 %arg, -9223372036854775808
+ %2 = bitcast i64 %1 to double
+ ret double %2
+}
+
+define double @bitcast_to_double_nan(i64 %arg) {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
+; CHECK-LABEL: define nofpclass(inf zero sub norm) double
@bitcast_to_double_nan
+; CHECK-SAME: (i64 [[ARG:%.*]]) #[[ATTR3]] {
+; CHECK-NEXT: [[TMP1:%.*]] = or i64 [[ARG]], -4503599627370495
+; CHECK-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to double
+; CHECK-NEXT: ret double [[TMP2]]
+;
+ %1 = or i64 %arg, -4503599627370495
+ %2 = bitcast i64 %1 to double
+ ret double %2
+}
+
declare i64 @_Z13get_global_idj(i32 noundef)
attributes #0 = { "denormal-fp-math"="preserve-sign,preserve-sign" }
>From 3db13768fb40b7a1dbe6bbd32dfeb9ed5c9215ed Mon Sep 17 00:00:00 2001
From: Alex MacLean <a...@alex-maclean.com>
Date: Fri, 5 Jul 2024 09:01:33 -0700
Subject: [PATCH 2/4] Update llvm/lib/Analysis/ValueTracking.cpp
Co-authored-by: Yingwei Zheng <dtcx...@qq.com>
---
llvm/lib/Analysis/ValueTracking.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/lib/Analysis/ValueTracking.cpp
b/llvm/lib/Analysis/ValueTracking.cpp
index a16c8e3d48403..35baffedbf255 100644
--- a/llvm/lib/Analysis/ValueTracking.cpp
+++ b/llvm/lib/Analysis/ValueTracking.cpp
@@ -5815,7 +5815,7 @@ void computeKnownFPClass(const Value *V, const APInt
&DemandedElts,
computeKnownBits(Casted, Bits, Depth + 1, Q);
// Transfer information from the sign bit.
- if (Bits.Zero.isSignBitSet())
+ if (Bits.isNonNegative())
Known.signBitMustBeZero();
else if (Bits.One.isSignBitSet())
Known.signBitMustBeOne();
>From 9061c04e98d61d28741ff1851be0892a5cf15227 Mon Sep 17 00:00:00 2001
From: Alex MacLean <a...@alex-maclean.com>
Date: Fri, 5 Jul 2024 09:01:49 -0700
Subject: [PATCH 3/4] Update llvm/lib/Analysis/ValueTracking.cpp
Co-authored-by: Yingwei Zheng <dtcx...@qq.com>
---
llvm/lib/Analysis/ValueTracking.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/lib/Analysis/ValueTracking.cpp
b/llvm/lib/Analysis/ValueTracking.cpp
index 35baffedbf255..ddbab998a097e 100644
--- a/llvm/lib/Analysis/ValueTracking.cpp
+++ b/llvm/lib/Analysis/ValueTracking.cpp
@@ -5817,7 +5817,7 @@ void computeKnownFPClass(const Value *V, const APInt
&DemandedElts,
// Transfer information from the sign bit.
if (Bits.isNonNegative())
Known.signBitMustBeZero();
- else if (Bits.One.isSignBitSet())
+ else if (Bits.isNegative())
Known.signBitMustBeOne();
if (Ty->isIEEE()) {
>From 0cc0f04e63fa58581bcb70d9443954a6b8727e15 Mon Sep 17 00:00:00 2001
From: Alex MacLean <amacl...@nvidia.com>
Date: Fri, 5 Jul 2024 17:12:22 +0000
Subject: [PATCH 4/4] address comments
---
clang/test/Headers/__clang_hip_math.hip | 571 ++++++++++++------
llvm/lib/Analysis/ValueTracking.cpp | 9 +-
.../CodeGen/AMDGPU/fneg-modifier-casting.ll | 20 -
llvm/test/Transforms/Attributor/nofpclass.ll | 149 +++--
4 files changed, 490 insertions(+), 259 deletions(-)
diff --git a/clang/test/Headers/__clang_hip_math.hip
b/clang/test/Headers/__clang_hip_math.hip
index 26da82843c512..0083bb0072468 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -2360,198 +2360,395 @@ extern "C" __device__ double test_modf(double x,
double* y) {
return modf(x, y);
}
-// CHECK-LABEL: @test_nanf(
-// CHECK-NEXT: entry:
-// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa
[[TBAA4]]
-// CHECK-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label
[[WHILE_COND_I14_I_I:%.*]]
-// CHECK: if.then.i.i:
-// CHECK-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr
[[TAG]], i64 1
-// CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1,
!tbaa [[TBAA4]]
-// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [
-// CHECK-NEXT: i8 120, label [[WHILE_COND_I30_I_I_PREHEADER:%.*]]
-// CHECK-NEXT: i8 88, label [[WHILE_COND_I30_I_I_PREHEADER]]
-// CHECK-NEXT: ]
-// CHECK: while.cond.i30.i.i.preheader:
-// CHECK-NEXT: br label [[WHILE_COND_I30_I_I:%.*]]
-// CHECK: while.cond.i30.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I37_I_I:%.*]], [[CLEANUP_I36_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]],
[[WHILE_COND_I30_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]],
[[CLEANUP_I36_I_I]] ], [ 0, [[WHILE_COND_I30_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align
1, !tbaa [[TBAA4]]
-// CHECK-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I33_I_I]], label [[_ZL4NANFPKC_EXIT:%.*]],
label [[WHILE_BODY_I34_I_I:%.*]]
-// CHECK: while.body.i34.i.i:
-// CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// CHECK-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I35_I_I]], label [[IF_END31_I_I_I:%.*]],
label [[IF_ELSE_I_I_I:%.*]]
-// CHECK: if.else.i.i.i:
-// CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
-// CHECK-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
-// CHECK-NEXT: br i1 [[OR_COND33_I_I_I]], label [[IF_END31_I_I_I]], label
[[IF_ELSE17_I_I_I:%.*]]
-// CHECK: if.else17.i.i.i:
-// CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
-// CHECK-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// CHECK-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label
[[CLEANUP_I36_I_I]]
-// CHECK: if.end31.i.i.i:
-// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I34_I_I]] ], [
-87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4
-// CHECK-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
-// CHECK-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
-// CHECK-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]],
[[CONV25_I_I_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8, ptr
[[__TAGP_ADDR_0_I31_I_I]], i64 1
-// CHECK-NEXT: br label [[CLEANUP_I36_I_I]]
-// CHECK: cleanup.i36.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [
[[INCDEC_PTR_I40_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]],
[[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]],
[[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [
false, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I30_I_I]], label
[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP11]]
-// CHECK: while.cond.i.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]],
[[IF_THEN_I_I]] ]
-// CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]],
[[CLEANUP_I_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
-// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align
1, !tbaa [[TBAA4]]
-// CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL4NANFPKC_EXIT]], label
[[WHILE_BODY_I_I_I:%.*]]
-// CHECK: while.body.i.i.i:
-// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// CHECK-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
-// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_THEN_I_I_I:%.*]], label
[[CLEANUP_I_I_I]]
-// CHECK: if.then.i.i.i:
-// CHECK-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
-// CHECK-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
-// CHECK-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
-// CHECK-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr
[[__TAGP_ADDR_0_I_I_I]], i64 1
-// CHECK-NEXT: br label [[CLEANUP_I_I_I]]
-// CHECK: cleanup.i.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]],
[[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
-// CHECK-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]], [[IF_THEN_I_I_I]]
], [ [[__R_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[WHILE_COND_I_I_I]], label
[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP7]]
-// CHECK: while.cond.i14.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I21_I_I:%.*]], [[CLEANUP_I20_I_I:%.*]] ], [ [[TAG]],
[[ENTRY:%.*]] ]
-// CHECK-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]],
[[CLEANUP_I20_I_I]] ], [ 0, [[ENTRY]] ]
-// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align
1, !tbaa [[TBAA4]]
-// CHECK-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I17_I_I]], label [[_ZL4NANFPKC_EXIT]], label
[[WHILE_BODY_I18_I_I:%.*]]
-// CHECK: while.body.i18.i.i:
-// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// CHECK-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I19_I_I]], label [[IF_THEN_I24_I_I:%.*]],
label [[CLEANUP_I20_I_I]]
-// CHECK: if.then.i24.i.i:
-// CHECK-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
-// CHECK-NEXT: [[CONV5_I26_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
-// CHECK-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48
-// CHECK-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]],
[[CONV5_I26_I_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds i8, ptr
[[__TAGP_ADDR_0_I15_I_I]], i64 1
-// CHECK-NEXT: br label [[CLEANUP_I20_I_I]]
-// CHECK: cleanup.i20.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [
[[INCDEC_PTR_I29_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]],
[[WHILE_BODY_I18_I_I]] ]
-// CHECK-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]],
[[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_BODY_I18_I_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I19_I_I]], label [[WHILE_COND_I14_I_I]],
label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP10]]
-// CHECK: _ZL4nanfPKc.exit:
-// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [
[[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I36_I_I]] ], [
[[__R_0_I32_I_I]], [[WHILE_COND_I30_I_I]] ], [ 0, [[CLEANUP_I20_I_I]] ], [
[[__R_0_I16_I_I]], [[WHILE_COND_I14_I_I]] ]
-// CHECK-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32
-// CHECK-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303
-// CHECK-NEXT: [[BF_SET9_I:%.*]] = or disjoint i32 [[BF_VALUE_I]],
2143289344
-// CHECK-NEXT: [[TMP10:%.*]] = bitcast i32 [[BF_SET9_I]] to float
-// CHECK-NEXT: ret float [[TMP10]]
+// DEFAULT-LABEL: @test_nanf(
+// DEFAULT-NEXT: entry:
+// DEFAULT-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa
[[TBAA4]]
+// DEFAULT-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
+// DEFAULT-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label
[[WHILE_COND_I14_I_I:%.*]]
+// DEFAULT: if.then.i.i:
+// DEFAULT-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr
[[TAG]], i64 1
+// DEFAULT-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1,
!tbaa [[TBAA4]]
+// DEFAULT-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [
+// DEFAULT-NEXT: i8 120, label [[WHILE_COND_I30_I_I_PREHEADER:%.*]]
+// DEFAULT-NEXT: i8 88, label [[WHILE_COND_I30_I_I_PREHEADER]]
+// DEFAULT-NEXT: ]
+// DEFAULT: while.cond.i30.i.i.preheader:
+// DEFAULT-NEXT: br label [[WHILE_COND_I30_I_I:%.*]]
+// DEFAULT: while.cond.i30.i.i:
+// DEFAULT-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I37_I_I:%.*]], [[CLEANUP_I36_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]],
[[WHILE_COND_I30_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]],
[[CLEANUP_I36_I_I]] ], [ 0, [[WHILE_COND_I30_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]],
align 1, !tbaa [[TBAA4]]
+// DEFAULT-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I33_I_I]], label [[_ZL4NANFPKC_EXIT:%.*]],
label [[WHILE_BODY_I34_I_I:%.*]]
+// DEFAULT: while.body.i34.i.i:
+// DEFAULT-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
+// DEFAULT-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
+// DEFAULT-NEXT: br i1 [[OR_COND_I35_I_I]], label [[IF_END31_I_I_I:%.*]],
label [[IF_ELSE_I_I_I:%.*]]
+// DEFAULT: if.else.i.i.i:
+// DEFAULT-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
+// DEFAULT-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// DEFAULT-NEXT: br i1 [[OR_COND33_I_I_I]], label [[IF_END31_I_I_I]], label
[[IF_ELSE17_I_I_I:%.*]]
+// DEFAULT: if.else17.i.i.i:
+// DEFAULT-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
+// DEFAULT-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
+// DEFAULT-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label
[[CLEANUP_I36_I_I]]
+// DEFAULT: if.end31.i.i.i:
+// DEFAULT-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I34_I_I]] ],
[ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
+// DEFAULT-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4
+// DEFAULT-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// DEFAULT-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
+// DEFAULT-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]],
[[CONV25_I_I_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8,
ptr [[__TAGP_ADDR_0_I31_I_I]], i64 1
+// DEFAULT-NEXT: br label [[CLEANUP_I36_I_I]]
+// DEFAULT: cleanup.i36.i.i:
+// DEFAULT-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [
[[INCDEC_PTR_I40_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]],
[[IF_ELSE17_I_I_I]] ]
+// DEFAULT-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]],
[[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], [[IF_ELSE17_I_I_I]] ]
+// DEFAULT-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ],
[ false, [[IF_ELSE17_I_I_I]] ]
+// DEFAULT-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I30_I_I]], label
[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP11]]
+// DEFAULT: while.cond.i.i.i:
+// DEFAULT-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]],
[[IF_THEN_I_I]] ]
+// DEFAULT-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]],
[[CLEANUP_I_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
+// DEFAULT-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align
1, !tbaa [[TBAA4]]
+// DEFAULT-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL4NANFPKC_EXIT]], label
[[WHILE_BODY_I_I_I:%.*]]
+// DEFAULT: while.body.i.i.i:
+// DEFAULT-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// DEFAULT-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// DEFAULT-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_THEN_I_I_I:%.*]],
label [[CLEANUP_I_I_I]]
+// DEFAULT: if.then.i.i.i:
+// DEFAULT-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
+// DEFAULT-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// DEFAULT-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
+// DEFAULT-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr
[[__TAGP_ADDR_0_I_I_I]], i64 1
+// DEFAULT-NEXT: br label [[CLEANUP_I_I_I]]
+// DEFAULT: cleanup.i.i.i:
+// DEFAULT-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]],
[[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// DEFAULT-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]],
[[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// DEFAULT-NEXT: br i1 [[OR_COND_I_I_I]], label [[WHILE_COND_I_I_I]], label
[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP7]]
+// DEFAULT: while.cond.i14.i.i:
+// DEFAULT-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I21_I_I:%.*]], [[CLEANUP_I20_I_I:%.*]] ], [ [[TAG]],
[[ENTRY:%.*]] ]
+// DEFAULT-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]],
[[CLEANUP_I20_I_I]] ], [ 0, [[ENTRY]] ]
+// DEFAULT-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]],
align 1, !tbaa [[TBAA4]]
+// DEFAULT-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I17_I_I]], label [[_ZL4NANFPKC_EXIT]],
label [[WHILE_BODY_I18_I_I:%.*]]
+// DEFAULT: while.body.i18.i.i:
+// DEFAULT-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// DEFAULT-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
+// DEFAULT-NEXT: br i1 [[OR_COND_I19_I_I]], label [[IF_THEN_I24_I_I:%.*]],
label [[CLEANUP_I20_I_I]]
+// DEFAULT: if.then.i24.i.i:
+// DEFAULT-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
+// DEFAULT-NEXT: [[CONV5_I26_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
+// DEFAULT-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48
+// DEFAULT-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]],
[[CONV5_I26_I_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds i8,
ptr [[__TAGP_ADDR_0_I15_I_I]], i64 1
+// DEFAULT-NEXT: br label [[CLEANUP_I20_I_I]]
+// DEFAULT: cleanup.i20.i.i:
+// DEFAULT-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [
[[INCDEC_PTR_I29_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]],
[[WHILE_BODY_I18_I_I]] ]
+// DEFAULT-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]],
[[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_BODY_I18_I_I]] ]
+// DEFAULT-NEXT: br i1 [[OR_COND_I19_I_I]], label [[WHILE_COND_I14_I_I]],
label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP10]]
+// DEFAULT: _ZL4nanfPKc.exit:
+// DEFAULT-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [
[[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I36_I_I]] ], [
[[__R_0_I32_I_I]], [[WHILE_COND_I30_I_I]] ], [ 0, [[CLEANUP_I20_I_I]] ], [
[[__R_0_I16_I_I]], [[WHILE_COND_I14_I_I]] ]
+// DEFAULT-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32
+// DEFAULT-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303
+// DEFAULT-NEXT: [[BF_SET9_I:%.*]] = or disjoint i32 [[BF_VALUE_I]],
2143289344
+// DEFAULT-NEXT: [[TMP10:%.*]] = bitcast i32 [[BF_SET9_I]] to float
+// DEFAULT-NEXT: ret float [[TMP10]]
+//
+// FINITEONLY-LABEL: @test_nanf(
+// FINITEONLY-NEXT: entry:
+// FINITEONLY-NEXT: ret float poison
+//
+// APPROX-LABEL: @test_nanf(
+// APPROX-NEXT: entry:
+// APPROX-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa
[[TBAA4]]
+// APPROX-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
+// APPROX-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label
[[WHILE_COND_I14_I_I:%.*]]
+// APPROX: if.then.i.i:
+// APPROX-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr
[[TAG]], i64 1
+// APPROX-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1,
!tbaa [[TBAA4]]
+// APPROX-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [
+// APPROX-NEXT: i8 120, label [[WHILE_COND_I30_I_I_PREHEADER:%.*]]
+// APPROX-NEXT: i8 88, label [[WHILE_COND_I30_I_I_PREHEADER]]
+// APPROX-NEXT: ]
+// APPROX: while.cond.i30.i.i.preheader:
+// APPROX-NEXT: br label [[WHILE_COND_I30_I_I:%.*]]
+// APPROX: while.cond.i30.i.i:
+// APPROX-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I37_I_I:%.*]], [[CLEANUP_I36_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]],
[[WHILE_COND_I30_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]],
[[CLEANUP_I36_I_I]] ], [ 0, [[WHILE_COND_I30_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]],
align 1, !tbaa [[TBAA4]]
+// APPROX-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
+// APPROX-NEXT: br i1 [[CMP_NOT_I33_I_I]], label [[_ZL4NANFPKC_EXIT:%.*]],
label [[WHILE_BODY_I34_I_I:%.*]]
+// APPROX: while.body.i34.i.i:
+// APPROX-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
+// APPROX-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
+// APPROX-NEXT: br i1 [[OR_COND_I35_I_I]], label [[IF_END31_I_I_I:%.*]],
label [[IF_ELSE_I_I_I:%.*]]
+// APPROX: if.else.i.i.i:
+// APPROX-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
+// APPROX-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// APPROX-NEXT: br i1 [[OR_COND33_I_I_I]], label [[IF_END31_I_I_I]], label
[[IF_ELSE17_I_I_I:%.*]]
+// APPROX: if.else17.i.i.i:
+// APPROX-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
+// APPROX-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
+// APPROX-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label
[[CLEANUP_I36_I_I]]
+// APPROX: if.end31.i.i.i:
+// APPROX-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I34_I_I]] ],
[ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
+// APPROX-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4
+// APPROX-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// APPROX-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
+// APPROX-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]],
[[CONV25_I_I_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8, ptr
[[__TAGP_ADDR_0_I31_I_I]], i64 1
+// APPROX-NEXT: br label [[CLEANUP_I36_I_I]]
+// APPROX: cleanup.i36.i.i:
+// APPROX-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [
[[INCDEC_PTR_I40_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]],
[[IF_ELSE17_I_I_I]] ]
+// APPROX-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]],
[[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], [[IF_ELSE17_I_I_I]] ]
+// APPROX-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [
false, [[IF_ELSE17_I_I_I]] ]
+// APPROX-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I30_I_I]], label
[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP11]]
+// APPROX: while.cond.i.i.i:
+// APPROX-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]],
[[IF_THEN_I_I]] ]
+// APPROX-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]],
[[CLEANUP_I_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
+// APPROX-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align
1, !tbaa [[TBAA4]]
+// APPROX-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// APPROX-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL4NANFPKC_EXIT]], label
[[WHILE_BODY_I_I_I:%.*]]
+// APPROX: while.body.i.i.i:
+// APPROX-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// APPROX-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// APPROX-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_THEN_I_I_I:%.*]], label
[[CLEANUP_I_I_I]]
+// APPROX: if.then.i.i.i:
+// APPROX-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
+// APPROX-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// APPROX-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
+// APPROX-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr
[[__TAGP_ADDR_0_I_I_I]], i64 1
+// APPROX-NEXT: br label [[CLEANUP_I_I_I]]
+// APPROX: cleanup.i.i.i:
+// APPROX-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]],
[[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// APPROX-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]],
[[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// APPROX-NEXT: br i1 [[OR_COND_I_I_I]], label [[WHILE_COND_I_I_I]], label
[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP7]]
+// APPROX: while.cond.i14.i.i:
+// APPROX-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I21_I_I:%.*]], [[CLEANUP_I20_I_I:%.*]] ], [ [[TAG]],
[[ENTRY:%.*]] ]
+// APPROX-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]],
[[CLEANUP_I20_I_I]] ], [ 0, [[ENTRY]] ]
+// APPROX-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]],
align 1, !tbaa [[TBAA4]]
+// APPROX-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
+// APPROX-NEXT: br i1 [[CMP_NOT_I17_I_I]], label [[_ZL4NANFPKC_EXIT]],
label [[WHILE_BODY_I18_I_I:%.*]]
+// APPROX: while.body.i18.i.i:
+// APPROX-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// APPROX-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
+// APPROX-NEXT: br i1 [[OR_COND_I19_I_I]], label [[IF_THEN_I24_I_I:%.*]],
label [[CLEANUP_I20_I_I]]
+// APPROX: if.then.i24.i.i:
+// APPROX-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
+// APPROX-NEXT: [[CONV5_I26_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
+// APPROX-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48
+// APPROX-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]],
[[CONV5_I26_I_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds i8, ptr
[[__TAGP_ADDR_0_I15_I_I]], i64 1
+// APPROX-NEXT: br label [[CLEANUP_I20_I_I]]
+// APPROX: cleanup.i20.i.i:
+// APPROX-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [
[[INCDEC_PTR_I29_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]],
[[WHILE_BODY_I18_I_I]] ]
+// APPROX-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]],
[[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_BODY_I18_I_I]] ]
+// APPROX-NEXT: br i1 [[OR_COND_I19_I_I]], label [[WHILE_COND_I14_I_I]],
label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP10]]
+// APPROX: _ZL4nanfPKc.exit:
+// APPROX-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [
[[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I36_I_I]] ], [
[[__R_0_I32_I_I]], [[WHILE_COND_I30_I_I]] ], [ 0, [[CLEANUP_I20_I_I]] ], [
[[__R_0_I16_I_I]], [[WHILE_COND_I14_I_I]] ]
+// APPROX-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32
+// APPROX-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303
+// APPROX-NEXT: [[BF_SET9_I:%.*]] = or disjoint i32 [[BF_VALUE_I]],
2143289344
+// APPROX-NEXT: [[TMP10:%.*]] = bitcast i32 [[BF_SET9_I]] to float
+// APPROX-NEXT: ret float [[TMP10]]
//
extern "C" __device__ float test_nanf(const char *tag) {
return nanf(tag);
}
-// CHECK-LABEL: @test_nan(
-// CHECK-NEXT: entry:
-// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa
[[TBAA4]]
-// CHECK-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label
[[WHILE_COND_I14_I_I:%.*]]
-// CHECK: if.then.i.i:
-// CHECK-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr
[[TAG]], i64 1
-// CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1,
!tbaa [[TBAA4]]
-// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [
-// CHECK-NEXT: i8 120, label [[WHILE_COND_I30_I_I_PREHEADER:%.*]]
-// CHECK-NEXT: i8 88, label [[WHILE_COND_I30_I_I_PREHEADER]]
-// CHECK-NEXT: ]
-// CHECK: while.cond.i30.i.i.preheader:
-// CHECK-NEXT: br label [[WHILE_COND_I30_I_I:%.*]]
-// CHECK: while.cond.i30.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I37_I_I:%.*]], [[CLEANUP_I36_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]],
[[WHILE_COND_I30_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]],
[[CLEANUP_I36_I_I]] ], [ 0, [[WHILE_COND_I30_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align
1, !tbaa [[TBAA4]]
-// CHECK-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I33_I_I]], label [[_ZL3NANPKC_EXIT:%.*]],
label [[WHILE_BODY_I34_I_I:%.*]]
-// CHECK: while.body.i34.i.i:
-// CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// CHECK-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I35_I_I]], label [[IF_END31_I_I_I:%.*]],
label [[IF_ELSE_I_I_I:%.*]]
-// CHECK: if.else.i.i.i:
-// CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
-// CHECK-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
-// CHECK-NEXT: br i1 [[OR_COND33_I_I_I]], label [[IF_END31_I_I_I]], label
[[IF_ELSE17_I_I_I:%.*]]
-// CHECK: if.else17.i.i.i:
-// CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
-// CHECK-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// CHECK-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label
[[CLEANUP_I36_I_I]]
-// CHECK: if.end31.i.i.i:
-// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I34_I_I]] ], [
-87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4
-// CHECK-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
-// CHECK-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
-// CHECK-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]],
[[CONV25_I_I_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8, ptr
[[__TAGP_ADDR_0_I31_I_I]], i64 1
-// CHECK-NEXT: br label [[CLEANUP_I36_I_I]]
-// CHECK: cleanup.i36.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [
[[INCDEC_PTR_I40_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]],
[[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]],
[[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [
false, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I30_I_I]], label
[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP11]]
-// CHECK: while.cond.i.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]],
[[IF_THEN_I_I]] ]
-// CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]],
[[CLEANUP_I_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
-// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align
1, !tbaa [[TBAA4]]
-// CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL3NANPKC_EXIT]], label
[[WHILE_BODY_I_I_I:%.*]]
-// CHECK: while.body.i.i.i:
-// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// CHECK-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
-// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_THEN_I_I_I:%.*]], label
[[CLEANUP_I_I_I]]
-// CHECK: if.then.i.i.i:
-// CHECK-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
-// CHECK-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
-// CHECK-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
-// CHECK-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr
[[__TAGP_ADDR_0_I_I_I]], i64 1
-// CHECK-NEXT: br label [[CLEANUP_I_I_I]]
-// CHECK: cleanup.i.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]],
[[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
-// CHECK-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]], [[IF_THEN_I_I_I]]
], [ [[__R_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[WHILE_COND_I_I_I]], label
[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP7]]
-// CHECK: while.cond.i14.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I21_I_I:%.*]], [[CLEANUP_I20_I_I:%.*]] ], [ [[TAG]],
[[ENTRY:%.*]] ]
-// CHECK-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]],
[[CLEANUP_I20_I_I]] ], [ 0, [[ENTRY]] ]
-// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align
1, !tbaa [[TBAA4]]
-// CHECK-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I17_I_I]], label [[_ZL3NANPKC_EXIT]], label
[[WHILE_BODY_I18_I_I:%.*]]
-// CHECK: while.body.i18.i.i:
-// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// CHECK-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I19_I_I]], label [[IF_THEN_I24_I_I:%.*]],
label [[CLEANUP_I20_I_I]]
-// CHECK: if.then.i24.i.i:
-// CHECK-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
-// CHECK-NEXT: [[CONV5_I26_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
-// CHECK-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48
-// CHECK-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]],
[[CONV5_I26_I_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds i8, ptr
[[__TAGP_ADDR_0_I15_I_I]], i64 1
-// CHECK-NEXT: br label [[CLEANUP_I20_I_I]]
-// CHECK: cleanup.i20.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [
[[INCDEC_PTR_I29_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]],
[[WHILE_BODY_I18_I_I]] ]
-// CHECK-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]],
[[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_BODY_I18_I_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I19_I_I]], label [[WHILE_COND_I14_I_I]],
label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP10]]
-// CHECK: _ZL3nanPKc.exit:
-// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [
[[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I36_I_I]] ], [
[[__R_0_I32_I_I]], [[WHILE_COND_I30_I_I]] ], [ 0, [[CLEANUP_I20_I_I]] ], [
[[__R_0_I16_I_I]], [[WHILE_COND_I14_I_I]] ]
-// CHECK-NEXT: [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]],
2251799813685247
-// CHECK-NEXT: [[BF_SET9_I:%.*]] = or disjoint i64 [[BF_VALUE_I]],
9221120237041090560
-// CHECK-NEXT: [[TMP10:%.*]] = bitcast i64 [[BF_SET9_I]] to double
-// CHECK-NEXT: ret double [[TMP10]]
+// DEFAULT-LABEL: @test_nan(
+// DEFAULT-NEXT: entry:
+// DEFAULT-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa
[[TBAA4]]
+// DEFAULT-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
+// DEFAULT-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label
[[WHILE_COND_I14_I_I:%.*]]
+// DEFAULT: if.then.i.i:
+// DEFAULT-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr
[[TAG]], i64 1
+// DEFAULT-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1,
!tbaa [[TBAA4]]
+// DEFAULT-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [
+// DEFAULT-NEXT: i8 120, label [[WHILE_COND_I30_I_I_PREHEADER:%.*]]
+// DEFAULT-NEXT: i8 88, label [[WHILE_COND_I30_I_I_PREHEADER]]
+// DEFAULT-NEXT: ]
+// DEFAULT: while.cond.i30.i.i.preheader:
+// DEFAULT-NEXT: br label [[WHILE_COND_I30_I_I:%.*]]
+// DEFAULT: while.cond.i30.i.i:
+// DEFAULT-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I37_I_I:%.*]], [[CLEANUP_I36_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]],
[[WHILE_COND_I30_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]],
[[CLEANUP_I36_I_I]] ], [ 0, [[WHILE_COND_I30_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]],
align 1, !tbaa [[TBAA4]]
+// DEFAULT-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I33_I_I]], label [[_ZL3NANPKC_EXIT:%.*]],
label [[WHILE_BODY_I34_I_I:%.*]]
+// DEFAULT: while.body.i34.i.i:
+// DEFAULT-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
+// DEFAULT-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
+// DEFAULT-NEXT: br i1 [[OR_COND_I35_I_I]], label [[IF_END31_I_I_I:%.*]],
label [[IF_ELSE_I_I_I:%.*]]
+// DEFAULT: if.else.i.i.i:
+// DEFAULT-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
+// DEFAULT-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// DEFAULT-NEXT: br i1 [[OR_COND33_I_I_I]], label [[IF_END31_I_I_I]], label
[[IF_ELSE17_I_I_I:%.*]]
+// DEFAULT: if.else17.i.i.i:
+// DEFAULT-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
+// DEFAULT-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
+// DEFAULT-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label
[[CLEANUP_I36_I_I]]
+// DEFAULT: if.end31.i.i.i:
+// DEFAULT-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I34_I_I]] ],
[ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
+// DEFAULT-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4
+// DEFAULT-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// DEFAULT-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
+// DEFAULT-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]],
[[CONV25_I_I_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8,
ptr [[__TAGP_ADDR_0_I31_I_I]], i64 1
+// DEFAULT-NEXT: br label [[CLEANUP_I36_I_I]]
+// DEFAULT: cleanup.i36.i.i:
+// DEFAULT-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [
[[INCDEC_PTR_I40_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]],
[[IF_ELSE17_I_I_I]] ]
+// DEFAULT-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]],
[[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], [[IF_ELSE17_I_I_I]] ]
+// DEFAULT-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ],
[ false, [[IF_ELSE17_I_I_I]] ]
+// DEFAULT-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I30_I_I]], label
[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP11]]
+// DEFAULT: while.cond.i.i.i:
+// DEFAULT-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]],
[[IF_THEN_I_I]] ]
+// DEFAULT-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]],
[[CLEANUP_I_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
+// DEFAULT-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align
1, !tbaa [[TBAA4]]
+// DEFAULT-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL3NANPKC_EXIT]], label
[[WHILE_BODY_I_I_I:%.*]]
+// DEFAULT: while.body.i.i.i:
+// DEFAULT-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// DEFAULT-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// DEFAULT-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_THEN_I_I_I:%.*]],
label [[CLEANUP_I_I_I]]
+// DEFAULT: if.then.i.i.i:
+// DEFAULT-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
+// DEFAULT-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// DEFAULT-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
+// DEFAULT-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr
[[__TAGP_ADDR_0_I_I_I]], i64 1
+// DEFAULT-NEXT: br label [[CLEANUP_I_I_I]]
+// DEFAULT: cleanup.i.i.i:
+// DEFAULT-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]],
[[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// DEFAULT-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]],
[[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// DEFAULT-NEXT: br i1 [[OR_COND_I_I_I]], label [[WHILE_COND_I_I_I]], label
[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP7]]
+// DEFAULT: while.cond.i14.i.i:
+// DEFAULT-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I21_I_I:%.*]], [[CLEANUP_I20_I_I:%.*]] ], [ [[TAG]],
[[ENTRY:%.*]] ]
+// DEFAULT-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]],
[[CLEANUP_I20_I_I]] ], [ 0, [[ENTRY]] ]
+// DEFAULT-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]],
align 1, !tbaa [[TBAA4]]
+// DEFAULT-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I17_I_I]], label [[_ZL3NANPKC_EXIT]],
label [[WHILE_BODY_I18_I_I:%.*]]
+// DEFAULT: while.body.i18.i.i:
+// DEFAULT-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// DEFAULT-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
+// DEFAULT-NEXT: br i1 [[OR_COND_I19_I_I]], label [[IF_THEN_I24_I_I:%.*]],
label [[CLEANUP_I20_I_I]]
+// DEFAULT: if.then.i24.i.i:
+// DEFAULT-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
+// DEFAULT-NEXT: [[CONV5_I26_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
+// DEFAULT-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48
+// DEFAULT-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]],
[[CONV5_I26_I_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds i8,
ptr [[__TAGP_ADDR_0_I15_I_I]], i64 1
+// DEFAULT-NEXT: br label [[CLEANUP_I20_I_I]]
+// DEFAULT: cleanup.i20.i.i:
+// DEFAULT-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [
[[INCDEC_PTR_I29_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]],
[[WHILE_BODY_I18_I_I]] ]
+// DEFAULT-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]],
[[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_BODY_I18_I_I]] ]
+// DEFAULT-NEXT: br i1 [[OR_COND_I19_I_I]], label [[WHILE_COND_I14_I_I]],
label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP10]]
+// DEFAULT: _ZL3nanPKc.exit:
+// DEFAULT-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [
[[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I36_I_I]] ], [
[[__R_0_I32_I_I]], [[WHILE_COND_I30_I_I]] ], [ 0, [[CLEANUP_I20_I_I]] ], [
[[__R_0_I16_I_I]], [[WHILE_COND_I14_I_I]] ]
+// DEFAULT-NEXT: [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]],
2251799813685247
+// DEFAULT-NEXT: [[BF_SET9_I:%.*]] = or disjoint i64 [[BF_VALUE_I]],
9221120237041090560
+// DEFAULT-NEXT: [[TMP10:%.*]] = bitcast i64 [[BF_SET9_I]] to double
+// DEFAULT-NEXT: ret double [[TMP10]]
+//
+// FINITEONLY-LABEL: @test_nan(
+// FINITEONLY-NEXT: entry:
+// FINITEONLY-NEXT: ret double poison
+//
+// APPROX-LABEL: @test_nan(
+// APPROX-NEXT: entry:
+// APPROX-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa
[[TBAA4]]
+// APPROX-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
+// APPROX-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label
[[WHILE_COND_I14_I_I:%.*]]
+// APPROX: if.then.i.i:
+// APPROX-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr
[[TAG]], i64 1
+// APPROX-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1,
!tbaa [[TBAA4]]
+// APPROX-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [
+// APPROX-NEXT: i8 120, label [[WHILE_COND_I30_I_I_PREHEADER:%.*]]
+// APPROX-NEXT: i8 88, label [[WHILE_COND_I30_I_I_PREHEADER]]
+// APPROX-NEXT: ]
+// APPROX: while.cond.i30.i.i.preheader:
+// APPROX-NEXT: br label [[WHILE_COND_I30_I_I:%.*]]
+// APPROX: while.cond.i30.i.i:
+// APPROX-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I37_I_I:%.*]], [[CLEANUP_I36_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]],
[[WHILE_COND_I30_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]],
[[CLEANUP_I36_I_I]] ], [ 0, [[WHILE_COND_I30_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]],
align 1, !tbaa [[TBAA4]]
+// APPROX-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
+// APPROX-NEXT: br i1 [[CMP_NOT_I33_I_I]], label [[_ZL3NANPKC_EXIT:%.*]],
label [[WHILE_BODY_I34_I_I:%.*]]
+// APPROX: while.body.i34.i.i:
+// APPROX-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
+// APPROX-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
+// APPROX-NEXT: br i1 [[OR_COND_I35_I_I]], label [[IF_END31_I_I_I:%.*]],
label [[IF_ELSE_I_I_I:%.*]]
+// APPROX: if.else.i.i.i:
+// APPROX-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
+// APPROX-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// APPROX-NEXT: br i1 [[OR_COND33_I_I_I]], label [[IF_END31_I_I_I]], label
[[IF_ELSE17_I_I_I:%.*]]
+// APPROX: if.else17.i.i.i:
+// APPROX-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
+// APPROX-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
+// APPROX-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label
[[CLEANUP_I36_I_I]]
+// APPROX: if.end31.i.i.i:
+// APPROX-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I34_I_I]] ],
[ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
+// APPROX-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4
+// APPROX-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// APPROX-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
+// APPROX-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]],
[[CONV25_I_I_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8, ptr
[[__TAGP_ADDR_0_I31_I_I]], i64 1
+// APPROX-NEXT: br label [[CLEANUP_I36_I_I]]
+// APPROX: cleanup.i36.i.i:
+// APPROX-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [
[[INCDEC_PTR_I40_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]],
[[IF_ELSE17_I_I_I]] ]
+// APPROX-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]],
[[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], [[IF_ELSE17_I_I_I]] ]
+// APPROX-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [
false, [[IF_ELSE17_I_I_I]] ]
+// APPROX-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I30_I_I]], label
[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP11]]
+// APPROX: while.cond.i.i.i:
+// APPROX-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]],
[[IF_THEN_I_I]] ]
+// APPROX-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]],
[[CLEANUP_I_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
+// APPROX-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align
1, !tbaa [[TBAA4]]
+// APPROX-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// APPROX-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL3NANPKC_EXIT]], label
[[WHILE_BODY_I_I_I:%.*]]
+// APPROX: while.body.i.i.i:
+// APPROX-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// APPROX-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// APPROX-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_THEN_I_I_I:%.*]], label
[[CLEANUP_I_I_I]]
+// APPROX: if.then.i.i.i:
+// APPROX-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
+// APPROX-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// APPROX-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
+// APPROX-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr
[[__TAGP_ADDR_0_I_I_I]], i64 1
+// APPROX-NEXT: br label [[CLEANUP_I_I_I]]
+// APPROX: cleanup.i.i.i:
+// APPROX-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]],
[[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// APPROX-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]],
[[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// APPROX-NEXT: br i1 [[OR_COND_I_I_I]], label [[WHILE_COND_I_I_I]], label
[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP7]]
+// APPROX: while.cond.i14.i.i:
+// APPROX-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [
[[__TAGP_ADDR_1_I21_I_I:%.*]], [[CLEANUP_I20_I_I:%.*]] ], [ [[TAG]],
[[ENTRY:%.*]] ]
+// APPROX-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]],
[[CLEANUP_I20_I_I]] ], [ 0, [[ENTRY]] ]
+// APPROX-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]],
align 1, !tbaa [[TBAA4]]
+// APPROX-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
+// APPROX-NEXT: br i1 [[CMP_NOT_I17_I_I]], label [[_ZL3NANPKC_EXIT]], label
[[WHILE_BODY_I18_I_I:%.*]]
+// APPROX: while.body.i18.i.i:
+// APPROX-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// APPROX-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
+// APPROX-NEXT: br i1 [[OR_COND_I19_I_I]], label [[IF_THEN_I24_I_I:%.*]],
label [[CLEANUP_I20_I_I]]
+// APPROX: if.then.i24.i.i:
+// APPROX-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
+// APPROX-NEXT: [[CONV5_I26_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
+// APPROX-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48
+// APPROX-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]],
[[CONV5_I26_I_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds i8, ptr
[[__TAGP_ADDR_0_I15_I_I]], i64 1
+// APPROX-NEXT: br label [[CLEANUP_I20_I_I]]
+// APPROX: cleanup.i20.i.i:
+// APPROX-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [
[[INCDEC_PTR_I29_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]],
[[WHILE_BODY_I18_I_I]] ]
+// APPROX-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]],
[[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_BODY_I18_I_I]] ]
+// APPROX-NEXT: br i1 [[OR_COND_I19_I_I]], label [[WHILE_COND_I14_I_I]],
label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP10]]
+// APPROX: _ZL3nanPKc.exit:
+// APPROX-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [
[[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I36_I_I]] ], [
[[__R_0_I32_I_I]], [[WHILE_COND_I30_I_I]] ], [ 0, [[CLEANUP_I20_I_I]] ], [
[[__R_0_I16_I_I]], [[WHILE_COND_I14_I_I]] ]
+// APPROX-NEXT: [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]],
2251799813685247
+// APPROX-NEXT: [[BF_SET9_I:%.*]] = or disjoint i64 [[BF_VALUE_I]],
9221120237041090560
+// APPROX-NEXT: [[TMP10:%.*]] = bitcast i64 [[BF_SET9_I]] to double
+// APPROX-NEXT: ret double [[TMP10]]
//
extern "C" __device__ double test_nan(const char *tag) {
return nan(tag);
diff --git a/llvm/lib/Analysis/ValueTracking.cpp
b/llvm/lib/Analysis/ValueTracking.cpp
index ddbab998a097e..8014592d82bbc 100644
--- a/llvm/lib/Analysis/ValueTracking.cpp
+++ b/llvm/lib/Analysis/ValueTracking.cpp
@@ -5806,13 +5806,14 @@ void computeKnownFPClass(const Value *V, const APInt
&DemandedElts,
break;
}
case Instruction::BitCast: {
- const Type *Ty = Op->getType();
- const Value *Casted = Op->getOperand(0);
- if (Ty->isVectorTy() || !Casted->getType()->isIntOrIntVectorTy())
+ const Value *Src;
+ if (!match(Op, m_ElementWiseBitCast(m_Value(Src))) ||
+ !Src->getType()->isIntOrIntVectorTy())
break;
+ const Type *Ty = Op->getType()->getScalarType();
KnownBits Bits(Ty->getScalarSizeInBits());
- computeKnownBits(Casted, Bits, Depth + 1, Q);
+ computeKnownBits(Src, DemandedElts, Bits, Depth + 1, Q);
// Transfer information from the sign bit.
if (Bits.isNonNegative())
diff --git a/llvm/test/CodeGen/AMDGPU/fneg-modifier-casting.ll
b/llvm/test/CodeGen/AMDGPU/fneg-modifier-casting.ll
index cd1ec85eb8d0f..df6ff7df029b4 100644
--- a/llvm/test/CodeGen/AMDGPU/fneg-modifier-casting.ll
+++ b/llvm/test/CodeGen/AMDGPU/fneg-modifier-casting.ll
@@ -1549,30 +1549,10 @@ define amdgpu_kernel void
@multiple_uses_fneg_select_f64(double %x, double %y, i
define amdgpu_kernel void @fnge_select_f32_multi_use_regression(float %.i2369)
{
; GCN-LABEL: fnge_select_f32_multi_use_regression:
; GCN: ; %bb.0: ; %.entry
-; GCN-NEXT: s_load_dword s0, s[4:5], 0x0
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: v_cmp_nlt_f32_e64 s[0:1], s0, 0
-; GCN-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[0:1]
-; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, 0, v0
-; GCN-NEXT: v_cndmask_b32_e32 v1, 0, v0, vcc
-; GCN-NEXT: v_mul_f32_e64 v0, -v0, v1
-; GCN-NEXT: v_cmp_lt_f32_e32 vcc, 0, v0
-; GCN-NEXT: s_and_b64 vcc, exec, vcc
; GCN-NEXT: s_endpgm
;
; GFX11-LABEL: fnge_select_f32_multi_use_regression:
; GFX11: ; %bb.0: ; %.entry
-; GFX11-NEXT: s_load_b32 s0, s[0:1], 0x0
-; GFX11-NEXT: s_waitcnt lgkmcnt(0)
-; GFX11-NEXT: v_cmp_nlt_f32_e64 s0, s0, 0
-; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) |
instid1(VALU_DEP_1)
-; GFX11-NEXT: v_cndmask_b32_e64 v0, 0, 1, s0
-; GFX11-NEXT: v_cmp_ngt_f32_e32 vcc_lo, 0, v0
-; GFX11-NEXT: v_cndmask_b32_e32 v1, 0, v0, vcc_lo
-; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) |
instid1(VALU_DEP_1)
-; GFX11-NEXT: v_mul_f32_e64 v0, -v0, v1
-; GFX11-NEXT: v_cmp_lt_f32_e32 vcc_lo, 0, v0
-; GFX11-NEXT: s_and_b32 vcc_lo, exec_lo, vcc_lo
; GFX11-NEXT: s_endpgm
.entry:
%i = fcmp uge float %.i2369, 0.000000e+00
diff --git a/llvm/test/Transforms/Attributor/nofpclass.ll
b/llvm/test/Transforms/Attributor/nofpclass.ll
index c5d562a436b33..5ea688113eb19 100644
--- a/llvm/test/Transforms/Attributor/nofpclass.ll
+++ b/llvm/test/Transforms/Attributor/nofpclass.ll
@@ -2694,104 +2694,157 @@ define float @bitcast_to_float_sign_0(i32 %arg) {
; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) float
@bitcast_to_float_sign_0
; CHECK-SAME: (i32 [[ARG:%.*]]) #[[ATTR3]] {
-; CHECK-NEXT: [[TMP1:%.*]] = lshr i32 [[ARG]], 1
-; CHECK-NEXT: [[TMP2:%.*]] = bitcast i32 [[TMP1]] to float
-; CHECK-NEXT: ret float [[TMP2]]
+; CHECK-NEXT: [[SHR:%.*]] = lshr i32 [[ARG]], 1
+; CHECK-NEXT: [[CAST:%.*]] = bitcast i32 [[SHR]] to float
+; CHECK-NEXT: ret float [[CAST]]
;
- %1 = lshr i32 %arg, 1
- %2 = bitcast i32 %1 to float
- ret float %2
+ %shr = lshr i32 %arg, 1
+ %cast = bitcast i32 %shr to float
+ ret float %cast
}
define float @bitcast_to_float_nnan(i32 %arg) {
; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
; CHECK-LABEL: define nofpclass(nan ninf nzero nsub nnorm) float
@bitcast_to_float_nnan
; CHECK-SAME: (i32 [[ARG:%.*]]) #[[ATTR3]] {
-; CHECK-NEXT: [[TMP1:%.*]] = lshr i32 [[ARG]], 2
-; CHECK-NEXT: [[TMP2:%.*]] = bitcast i32 [[TMP1]] to float
-; CHECK-NEXT: ret float [[TMP2]]
+; CHECK-NEXT: [[SHR:%.*]] = lshr i32 [[ARG]], 2
+; CHECK-NEXT: [[CAST:%.*]] = bitcast i32 [[SHR]] to float
+; CHECK-NEXT: ret float [[CAST]]
;
- %1 = lshr i32 %arg, 2
- %2 = bitcast i32 %1 to float
- ret float %2
+ %shr = lshr i32 %arg, 2
+ %cast = bitcast i32 %shr to float
+ ret float %cast
}
define float @bitcast_to_float_sign_1(i32 %arg) {
; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) float
@bitcast_to_float_sign_1
; CHECK-SAME: (i32 [[ARG:%.*]]) #[[ATTR3]] {
-; CHECK-NEXT: [[TMP1:%.*]] = or i32 [[ARG]], -2147483648
-; CHECK-NEXT: [[TMP2:%.*]] = bitcast i32 [[TMP1]] to float
-; CHECK-NEXT: ret float [[TMP2]]
+; CHECK-NEXT: [[OR:%.*]] = or i32 [[ARG]], -2147483648
+; CHECK-NEXT: [[CAST:%.*]] = bitcast i32 [[OR]] to float
+; CHECK-NEXT: ret float [[CAST]]
;
- %1 = or i32 %arg, -2147483648
- %2 = bitcast i32 %1 to float
- ret float %2
+ %or = or i32 %arg, -2147483648
+ %cast = bitcast i32 %or to float
+ ret float %cast
}
define float @bitcast_to_float_nan(i32 %arg) {
; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
; CHECK-LABEL: define nofpclass(inf zero sub norm) float @bitcast_to_float_nan
; CHECK-SAME: (i32 [[ARG:%.*]]) #[[ATTR3]] {
-; CHECK-NEXT: [[TMP1:%.*]] = or i32 [[ARG]], 2139095041
-; CHECK-NEXT: [[TMP2:%.*]] = bitcast i32 [[TMP1]] to float
-; CHECK-NEXT: ret float [[TMP2]]
+; CHECK-NEXT: [[OR:%.*]] = or i32 [[ARG]], 2139095041
+; CHECK-NEXT: [[CAST:%.*]] = bitcast i32 [[OR]] to float
+; CHECK-NEXT: ret float [[CAST]]
;
- %1 = or i32 %arg, 2139095041
- %2 = bitcast i32 %1 to float
- ret float %2
+ %or = or i32 %arg, 2139095041
+ %cast = bitcast i32 %or to float
+ ret float %cast
}
define double @bitcast_to_double_sign_0(i64 %arg) {
; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) double
@bitcast_to_double_sign_0
; CHECK-SAME: (i64 [[ARG:%.*]]) #[[ATTR3]] {
-; CHECK-NEXT: [[TMP1:%.*]] = lshr i64 [[ARG]], 1
-; CHECK-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to double
-; CHECK-NEXT: ret double [[TMP2]]
+; CHECK-NEXT: [[SHR:%.*]] = lshr i64 [[ARG]], 1
+; CHECK-NEXT: [[CAST:%.*]] = bitcast i64 [[SHR]] to double
+; CHECK-NEXT: ret double [[CAST]]
;
- %1 = lshr i64 %arg, 1
- %2 = bitcast i64 %1 to double
- ret double %2
+ %shr = lshr i64 %arg, 1
+ %cast = bitcast i64 %shr to double
+ ret double %cast
}
define double @bitcast_to_double_nnan(i64 %arg) {
; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
; CHECK-LABEL: define nofpclass(nan ninf nzero nsub nnorm) double
@bitcast_to_double_nnan
; CHECK-SAME: (i64 [[ARG:%.*]]) #[[ATTR3]] {
-; CHECK-NEXT: [[TMP1:%.*]] = lshr i64 [[ARG]], 2
-; CHECK-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to double
-; CHECK-NEXT: ret double [[TMP2]]
+; CHECK-NEXT: [[SHR:%.*]] = lshr i64 [[ARG]], 2
+; CHECK-NEXT: [[CAST:%.*]] = bitcast i64 [[SHR]] to double
+; CHECK-NEXT: ret double [[CAST]]
;
- %1 = lshr i64 %arg, 2
- %2 = bitcast i64 %1 to double
- ret double %2
+ %shr = lshr i64 %arg, 2
+ %cast = bitcast i64 %shr to double
+ ret double %cast
}
define double @bitcast_to_double_sign_1(i64 %arg) {
; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) double
@bitcast_to_double_sign_1
; CHECK-SAME: (i64 [[ARG:%.*]]) #[[ATTR3]] {
-; CHECK-NEXT: [[TMP1:%.*]] = or i64 [[ARG]], -9223372036854775808
-; CHECK-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to double
-; CHECK-NEXT: ret double [[TMP2]]
+; CHECK-NEXT: [[OR:%.*]] = or i64 [[ARG]], -9223372036854775808
+; CHECK-NEXT: [[CAST:%.*]] = bitcast i64 [[OR]] to double
+; CHECK-NEXT: ret double [[CAST]]
;
- %1 = or i64 %arg, -9223372036854775808
- %2 = bitcast i64 %1 to double
- ret double %2
+ %or = or i64 %arg, -9223372036854775808
+ %cast = bitcast i64 %or to double
+ ret double %cast
}
define double @bitcast_to_double_nan(i64 %arg) {
; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
; CHECK-LABEL: define nofpclass(inf zero sub norm) double
@bitcast_to_double_nan
; CHECK-SAME: (i64 [[ARG:%.*]]) #[[ATTR3]] {
-; CHECK-NEXT: [[TMP1:%.*]] = or i64 [[ARG]], -4503599627370495
-; CHECK-NEXT: [[TMP2:%.*]] = bitcast i64 [[TMP1]] to double
-; CHECK-NEXT: ret double [[TMP2]]
+; CHECK-NEXT: [[OR:%.*]] = or i64 [[ARG]], -4503599627370495
+; CHECK-NEXT: [[CAST:%.*]] = bitcast i64 [[OR]] to double
+; CHECK-NEXT: ret double [[CAST]]
;
- %1 = or i64 %arg, -4503599627370495
- %2 = bitcast i64 %1 to double
- ret double %2
+ %or = or i64 %arg, -4503599627370495
+ %cast = bitcast i64 %or to double
+ ret double %cast
+}
+
+
+define <2 x float> @bitcast_to_float_vect_sign_0(<2 x i32> %arg) {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
+; CHECK-LABEL: define nofpclass(ninf nzero nsub nnorm) <2 x float>
@bitcast_to_float_vect_sign_0
+; CHECK-SAME: (<2 x i32> [[ARG:%.*]]) #[[ATTR3]] {
+; CHECK-NEXT: [[SHR:%.*]] = lshr <2 x i32> [[ARG]], <i32 1, i32 2>
+; CHECK-NEXT: [[CAST:%.*]] = bitcast <2 x i32> [[SHR]] to <2 x float>
+; CHECK-NEXT: ret <2 x float> [[CAST]]
+;
+ %shr = lshr <2 x i32> %arg, <i32 1, i32 2>
+ %cast = bitcast <2 x i32> %shr to <2 x float>
+ ret <2 x float> %cast
+}
+
+define <2 x float> @bitcast_to_float_vect_nnan(<2 x i32> %arg) {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
+; CHECK-LABEL: define nofpclass(nan ninf nzero nsub nnorm) <2 x float>
@bitcast_to_float_vect_nnan
+; CHECK-SAME: (<2 x i32> [[ARG:%.*]]) #[[ATTR3]] {
+; CHECK-NEXT: [[SHR:%.*]] = lshr <2 x i32> [[ARG]], <i32 4, i32 4>
+; CHECK-NEXT: [[CAST:%.*]] = bitcast <2 x i32> [[SHR]] to <2 x float>
+; CHECK-NEXT: ret <2 x float> [[CAST]]
+;
+ %shr = lshr <2 x i32> %arg, <i32 4, i32 4>
+ %cast = bitcast <2 x i32> %shr to <2 x float>
+ ret <2 x float> %cast
+}
+
+define <2 x float> @bitcast_to_float_vect_sign_1(<2 x i32> %arg) {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
+; CHECK-LABEL: define nofpclass(pinf pzero psub pnorm) <2 x float>
@bitcast_to_float_vect_sign_1
+; CHECK-SAME: (<2 x i32> [[ARG:%.*]]) #[[ATTR3]] {
+; CHECK-NEXT: [[OR:%.*]] = or <2 x i32> [[ARG]], <i32 -2147483648, i32
-2147483648>
+; CHECK-NEXT: [[CAST:%.*]] = bitcast <2 x i32> [[OR]] to <2 x float>
+; CHECK-NEXT: ret <2 x float> [[CAST]]
+;
+ %or = or <2 x i32> %arg, <i32 -2147483648, i32 -2147483648>
+ %cast = bitcast <2 x i32> %or to <2 x float>
+ ret <2 x float> %cast
+}
+
+define <2 x float> @bitcast_to_float_vect_nan(<2 x i32> %arg) {
+; CHECK: Function Attrs: mustprogress nofree norecurse nosync nounwind
willreturn memory(none)
+; CHECK-LABEL: define nofpclass(inf zero sub norm) <2 x float>
@bitcast_to_float_vect_nan
+; CHECK-SAME: (<2 x i32> [[ARG:%.*]]) #[[ATTR3]] {
+; CHECK-NEXT: [[OR:%.*]] = or <2 x i32> [[ARG]], <i32 2139095041, i32
2139095041>
+; CHECK-NEXT: [[CAST:%.*]] = bitcast <2 x i32> [[OR]] to <2 x float>
+; CHECK-NEXT: ret <2 x float> [[CAST]]
+;
+ %or = or <2 x i32> %arg, <i32 2139095041, i32 2139095041>
+ %cast = bitcast <2 x i32> %or to <2 x float>
+ ret <2 x float> %cast
}
declare i64 @_Z13get_global_idj(i32 noundef)
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
