[clang] 1a2e77c - Revert "Revert "Inlining: Run the legacy AlwaysInliner before the regular inliner.""
Amara Emerson via cfe-commits
cfe-commits at lists.llvm.org
Sat Oct 28 23:21:50 PDT 2023
Author: Amara Emerson
Date: 2023-10-28T23:21:11-07:00
New Revision: 1a2e77cf9e11dbf56b5720c607313a566eebb16e
URL: https://github.com/llvm/llvm-project/commit/1a2e77cf9e11dbf56b5720c607313a566eebb16e
DIFF: https://github.com/llvm/llvm-project/commit/1a2e77cf9e11dbf56b5720c607313a566eebb16e.diff
LOG: Revert "Revert "Inlining: Run the legacy AlwaysInliner before the regular inliner.""
This reverts commit 86bfeb906e3a95ae428f3e97d78d3d22a7c839f3.
This is a long time coming re-application that was originally reverted due to
regressions, unrelated to the actual inlining change. These regressions have since
been fixed due to another long-in-the-making change of a66051c6 landing.
Original commit message for reference:
---
We have several situations where it's beneficial for code size to ensure that every
call to always-inline functions are inlined before normal inlining decisions are
made. While the normal inliner runs in a "MandatoryOnly" mode to try to do this,
it only does it on a per-SCC basis, rather than the whole module. Ensuring that
all mandatory inlinings are done before any heuristic based decisions are made
just makes sense.
Despite being referred to the "legacy" AlwaysInliner pass, it's already necessary
for -O0 because the CGSCC inliner is too expensive in compile time to run at -O0.
This also fixes an exponential compile time blow up in
https://github.com/llvm/llvm-project/issues/59126
Differential Revision: https://reviews.llvm.org/D143624
---
Added:
llvm/test/Transforms/Inline/always-inline-phase-ordering.ll
Modified:
clang/test/Frontend/optimization-remark-with-hotness-new-pm.c
clang/test/Headers/__clang_hip_math.hip
clang/test/OpenMP/bug57757.cpp
llvm/lib/Passes/PassBuilderPipelines.cpp
llvm/test/Other/new-pm-defaults.ll
llvm/test/Other/new-pm-print-pipeline.ll
llvm/test/Other/new-pm-thinlto-postlink-defaults.ll
llvm/test/Other/new-pm-thinlto-postlink-pgo-defaults.ll
llvm/test/Other/new-pm-thinlto-postlink-samplepgo-defaults.ll
llvm/test/Other/new-pm-thinlto-prelink-defaults.ll
llvm/test/Other/new-pm-thinlto-prelink-pgo-defaults.ll
llvm/test/Other/new-pm-thinlto-prelink-samplepgo-defaults.ll
Removed:
################################################################################
diff --git a/clang/test/Frontend/optimization-remark-with-hotness-new-pm.c b/clang/test/Frontend/optimization-remark-with-hotness-new-pm.c
index cb1992ec4cb8576..173b43ba41bfede 100644
--- a/clang/test/Frontend/optimization-remark-with-hotness-new-pm.c
+++ b/clang/test/Frontend/optimization-remark-with-hotness-new-pm.c
@@ -73,7 +73,7 @@ void bar(int x) {
// THRESHOLD-NOT: hotness
// NO_PGO: '-fdiagnostics-show-hotness' requires profile-guided optimization information
// NO_PGO: '-fdiagnostics-hotness-threshold=' requires profile-guided optimization information
- // expected-remark at +1 {{'foo' inlined into 'bar': always inline attribute at callsite bar:8:10; (hotness:}}
+ // expected-remark at +1 {{'foo' inlined into 'bar' with (cost=always): always inline attribute at callsite bar:8:10; (hotness:}}
sum += foo(x, x - 2);
}
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index fc18e14d8229635..b57c38d70b14c79 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -138,26 +138,26 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[P:%.*]], align 1, !tbaa [[TBAA4]]
// CHECK-NEXT: [[CMP_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// CHECK-NEXT: br i1 [[CMP_I]], label [[IF_THEN_I:%.*]], label [[WHILE_COND_I30_I:%.*]]
+// CHECK-NEXT: br i1 [[CMP_I]], label [[IF_THEN_I:%.*]], label [[WHILE_COND_I14_I:%.*]]
// CHECK: if.then.i:
// CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds i8, ptr [[P]], i64 1
// CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I]], align 1, !tbaa [[TBAA4]]
-// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I14_I:%.*]] [
-// CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_PREHEADER:%.*]]
-// CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_PREHEADER]]
+// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I:%.*]] [
+// CHECK-NEXT: i8 120, label [[WHILE_COND_I30_I_PREHEADER:%.*]]
+// CHECK-NEXT: i8 88, label [[WHILE_COND_I30_I_PREHEADER]]
// CHECK-NEXT: ]
-// CHECK: while.cond.i.i.preheader:
-// CHECK-NEXT: br label [[WHILE_COND_I_I:%.*]]
-// CHECK: while.cond.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I:%.*]], [[CLEANUP_I_I:%.*]] ], [ [[INCDEC_PTR_I]], [[WHILE_COND_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[__R_0_I_I:%.*]] = phi i64 [ [[__R_2_I_I:%.*]], [[CLEANUP_I_I]] ], [ 0, [[WHILE_COND_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I]], align 1, !tbaa [[TBAA4]]
-// CHECK-NEXT: [[CMP_NOT_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT:%.*]], label [[WHILE_BODY_I_I:%.*]]
-// CHECK: while.body.i.i:
+// CHECK: while.cond.i30.i.preheader:
+// CHECK-NEXT: br label [[WHILE_COND_I30_I:%.*]]
+// CHECK: while.cond.i30.i:
+// CHECK-NEXT: [[__TAGP_ADDR_0_I31_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I:%.*]], [[CLEANUP_I36_I:%.*]] ], [ [[INCDEC_PTR_I]], [[WHILE_COND_I30_I_PREHEADER]] ]
+// CHECK-NEXT: [[__R_0_I32_I:%.*]] = phi i64 [ [[__R_2_I_I:%.*]], [[CLEANUP_I36_I]] ], [ 0, [[WHILE_COND_I30_I_PREHEADER]] ]
+// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I]], align 1, !tbaa [[TBAA4]]
+// CHECK-NEXT: [[CMP_NOT_I33_I:%.*]] = icmp eq i8 [[TMP2]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I33_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT:%.*]], label [[WHILE_BODY_I34_I:%.*]]
+// CHECK: while.body.i34.i:
// CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// CHECK-NEXT: [[OR_COND_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I_I]], label [[IF_END31_I_I:%.*]], label [[IF_ELSE_I_I:%.*]]
+// CHECK-NEXT: [[OR_COND_I35_I:%.*]] = icmp ult i8 [[TMP3]], 10
+// CHECK-NEXT: br i1 [[OR_COND_I35_I]], label [[IF_END31_I_I:%.*]], label [[IF_ELSE_I_I:%.*]]
// CHECK: if.else.i.i:
// CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
// CHECK-NEXT: [[OR_COND33_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
@@ -165,33 +165,54 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
// CHECK: if.else17.i.i:
// CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
// CHECK-NEXT: [[OR_COND34_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// CHECK-NEXT: br i1 [[OR_COND34_I_I]], label [[IF_END31_I_I]], label [[CLEANUP_I_I]]
+// CHECK-NEXT: br i1 [[OR_COND34_I_I]], label [[IF_END31_I_I]], label [[CLEANUP_I36_I]]
// CHECK: if.end31.i.i:
-// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I]] ], [ -87, [[IF_ELSE_I_I]] ], [ -55, [[IF_ELSE17_I_I]] ]
-// CHECK-NEXT: [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I_I]], 4
+// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I34_I]] ], [ -87, [[IF_ELSE_I_I]] ], [ -55, [[IF_ELSE17_I_I]] ]
+// CHECK-NEXT: [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I32_I]], 4
// CHECK-NEXT: [[CONV25_I_I:%.*]] = sext i8 [[TMP2]] to i64
// CHECK-NEXT: [[ADD26_I_I:%.*]] = add i64 [[MUL24_I_I]], [[DOTSINK]]
// CHECK-NEXT: [[ADD28_I_I:%.*]] = add i64 [[ADD26_I_I]], [[CONV25_I_I]]
+// CHECK-NEXT: [[INCDEC_PTR_I40_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I31_I]], i64 1
+// CHECK-NEXT: br label [[CLEANUP_I36_I]]
+// CHECK: cleanup.i36.i:
+// CHECK-NEXT: [[__TAGP_ADDR_1_I37_I]] = phi ptr [ [[INCDEC_PTR_I40_I]], [[IF_END31_I_I]] ], [ [[__TAGP_ADDR_0_I31_I]], [[IF_ELSE17_I_I]] ]
+// CHECK-NEXT: [[__R_2_I_I]] = phi i64 [ [[ADD28_I_I]], [[IF_END31_I_I]] ], [ [[__R_0_I32_I]], [[IF_ELSE17_I_I]] ]
+// CHECK-NEXT: [[COND_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I]] ], [ false, [[IF_ELSE17_I_I]] ]
+// CHECK-NEXT: br i1 [[COND_I_I]], label [[WHILE_COND_I30_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP11]]
+// CHECK: while.cond.i.i:
+// CHECK-NEXT: [[__TAGP_ADDR_0_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I:%.*]], [[CLEANUP_I_I:%.*]] ], [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ]
+// CHECK-NEXT: [[__R_0_I_I:%.*]] = phi i64 [ [[__R_1_I_I:%.*]], [[CLEANUP_I_I]] ], [ 0, [[IF_THEN_I]] ]
+// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I]], align 1, !tbaa [[TBAA4]]
+// CHECK-NEXT: [[CMP_NOT_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I_I:%.*]]
+// CHECK: while.body.i.i:
+// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// CHECK-NEXT: [[OR_COND_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// CHECK-NEXT: br i1 [[OR_COND_I_I]], label [[IF_THEN_I_I:%.*]], label [[CLEANUP_I_I]]
+// CHECK: if.then.i.i:
+// CHECK-NEXT: [[MUL_I_I:%.*]] = shl i64 [[__R_0_I_I]], 3
+// CHECK-NEXT: [[CONV5_I_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT: [[ADD_I_I:%.*]] = add i64 [[MUL_I_I]], -48
+// CHECK-NEXT: [[SUB_I_I:%.*]] = add i64 [[ADD_I_I]], [[CONV5_I_I]]
// CHECK-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I]], i64 1
// CHECK-NEXT: br label [[CLEANUP_I_I]]
// CHECK: cleanup.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I_I]] = phi ptr [ [[INCDEC_PTR_I_I]], [[IF_END31_I_I]] ], [ [[__TAGP_ADDR_0_I_I]], [[IF_ELSE17_I_I]] ]
-// CHECK-NEXT: [[__R_2_I_I]] = phi i64 [ [[ADD28_I_I]], [[IF_END31_I_I]] ], [ [[__R_0_I_I]], [[IF_ELSE17_I_I]] ]
-// CHECK-NEXT: [[COND_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I]] ], [ false, [[IF_ELSE17_I_I]] ]
-// CHECK-NEXT: br i1 [[COND_I_I]], label [[WHILE_COND_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP11]]
+// CHECK-NEXT: [[__TAGP_ADDR_1_I_I]] = phi ptr [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ], [ [[__TAGP_ADDR_0_I_I]], [[WHILE_BODY_I_I]] ]
+// CHECK-NEXT: [[__R_1_I_I]] = phi i64 [ [[SUB_I_I]], [[IF_THEN_I_I]] ], [ [[__R_0_I_I]], [[WHILE_BODY_I_I]] ]
+// CHECK-NEXT: br i1 [[OR_COND_I_I]], label [[WHILE_COND_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP7]]
// CHECK: while.cond.i14.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I15_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I:%.*]], [[CLEANUP_I20_I:%.*]] ], [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ]
-// CHECK-NEXT: [[__R_0_I16_I:%.*]] = phi i64 [ [[__R_1_I22_I:%.*]], [[CLEANUP_I20_I]] ], [ 0, [[IF_THEN_I]] ]
-// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I]], align 1, !tbaa [[TBAA4]]
-// CHECK-NEXT: [[CMP_NOT_I17_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// CHECK-NEXT: [[__TAGP_ADDR_0_I15_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I:%.*]], [[CLEANUP_I20_I:%.*]] ], [ [[P]], [[ENTRY:%.*]] ]
+// CHECK-NEXT: [[__R_0_I16_I:%.*]] = phi i64 [ [[__R_1_I22_I:%.*]], [[CLEANUP_I20_I]] ], [ 0, [[ENTRY]] ]
+// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I]], align 1, !tbaa [[TBAA4]]
+// CHECK-NEXT: [[CMP_NOT_I17_I:%.*]] = icmp eq i8 [[TMP8]], 0
// CHECK-NEXT: br i1 [[CMP_NOT_I17_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I18_I:%.*]]
// CHECK: while.body.i18.i:
-// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// CHECK-NEXT: [[OR_COND_I19_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// CHECK-NEXT: [[OR_COND_I19_I:%.*]] = icmp ult i8 [[TMP9]], 10
// CHECK-NEXT: br i1 [[OR_COND_I19_I]], label [[IF_THEN_I24_I:%.*]], label [[CLEANUP_I20_I]]
// CHECK: if.then.i24.i:
-// CHECK-NEXT: [[MUL_I25_I:%.*]] = shl i64 [[__R_0_I16_I]], 3
-// CHECK-NEXT: [[CONV5_I26_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT: [[MUL_I25_I:%.*]] = mul i64 [[__R_0_I16_I]], 10
+// CHECK-NEXT: [[CONV5_I26_I:%.*]] = sext i8 [[TMP8]] to i64
// CHECK-NEXT: [[ADD_I27_I:%.*]] = add i64 [[MUL_I25_I]], -48
// CHECK-NEXT: [[SUB_I28_I:%.*]] = add i64 [[ADD_I27_I]], [[CONV5_I26_I]]
// CHECK-NEXT: [[INCDEC_PTR_I29_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I15_I]], i64 1
@@ -199,30 +220,9 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
// CHECK: cleanup.i20.i:
// CHECK-NEXT: [[__TAGP_ADDR_1_I21_I]] = phi ptr [ [[INCDEC_PTR_I29_I]], [[IF_THEN_I24_I]] ], [ [[__TAGP_ADDR_0_I15_I]], [[WHILE_BODY_I18_I]] ]
// CHECK-NEXT: [[__R_1_I22_I]] = phi i64 [ [[SUB_I28_I]], [[IF_THEN_I24_I]] ], [ [[__R_0_I16_I]], [[WHILE_BODY_I18_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I19_I]], label [[WHILE_COND_I14_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP7]]
-// CHECK: while.cond.i30.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I31_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I:%.*]], [[CLEANUP_I36_I:%.*]] ], [ [[P]], [[ENTRY:%.*]] ]
-// CHECK-NEXT: [[__R_0_I32_I:%.*]] = phi i64 [ [[__R_1_I38_I:%.*]], [[CLEANUP_I36_I]] ], [ 0, [[ENTRY]] ]
-// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I]], align 1, !tbaa [[TBAA4]]
-// CHECK-NEXT: [[CMP_NOT_I33_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I33_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I34_I:%.*]]
-// CHECK: while.body.i34.i:
-// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// CHECK-NEXT: [[OR_COND_I35_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I35_I]], label [[IF_THEN_I40_I:%.*]], label [[CLEANUP_I36_I]]
-// CHECK: if.then.i40.i:
-// CHECK-NEXT: [[MUL_I41_I:%.*]] = mul i64 [[__R_0_I32_I]], 10
-// CHECK-NEXT: [[CONV5_I42_I:%.*]] = sext i8 [[TMP8]] to i64
-// CHECK-NEXT: [[ADD_I43_I:%.*]] = add i64 [[MUL_I41_I]], -48
-// CHECK-NEXT: [[SUB_I44_I:%.*]] = add i64 [[ADD_I43_I]], [[CONV5_I42_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I45_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I31_I]], i64 1
-// CHECK-NEXT: br label [[CLEANUP_I36_I]]
-// CHECK: cleanup.i36.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I37_I]] = phi ptr [ [[INCDEC_PTR_I45_I]], [[IF_THEN_I40_I]] ], [ [[__TAGP_ADDR_0_I31_I]], [[WHILE_BODY_I34_I]] ]
-// CHECK-NEXT: [[__R_1_I38_I]] = phi i64 [ [[SUB_I44_I]], [[IF_THEN_I40_I]] ], [ [[__R_0_I32_I]], [[WHILE_BODY_I34_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I35_I]], label [[WHILE_COND_I30_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP10]]
+// CHECK-NEXT: br i1 [[OR_COND_I19_I]], label [[WHILE_COND_I14_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP10]]
// CHECK: _ZL15__make_mantissaPKc.exit:
-// CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, [[CLEANUP_I20_I]] ], [ [[__R_0_I16_I]], [[WHILE_COND_I14_I]] ], [ 0, [[CLEANUP_I_I]] ], [ [[__R_0_I_I]], [[WHILE_COND_I_I]] ], [ 0, [[CLEANUP_I36_I]] ], [ [[__R_0_I32_I]], [[WHILE_COND_I30_I]] ]
+// CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I]] ], [ [[__R_0_I_I]], [[WHILE_COND_I_I]] ], [ 0, [[CLEANUP_I36_I]] ], [ [[__R_0_I32_I]], [[WHILE_COND_I30_I]] ], [ 0, [[CLEANUP_I20_I]] ], [ [[__R_0_I16_I]], [[WHILE_COND_I14_I]] ]
// CHECK-NEXT: ret i64 [[RETVAL_0_I]]
//
extern "C" __device__ uint64_t test___make_mantissa(const char *p) {
@@ -231,8 +231,8 @@ extern "C" __device__ uint64_t test___make_mantissa(const char *p) {
// CHECK-LABEL: @test_abs(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[ABS_I:%.*]] = tail call noundef i32 @llvm.abs.i32(i32 [[X:%.*]], i1 true)
-// CHECK-NEXT: ret i32 [[ABS_I]]
+// CHECK-NEXT: [[TMP0:%.*]] = tail call noundef i32 @llvm.abs.i32(i32 [[X:%.*]], i1 true)
+// CHECK-NEXT: ret i32 [[TMP0]]
//
extern "C" __device__ int test_abs(int x) {
return abs(x);
@@ -240,8 +240,8 @@ extern "C" __device__ int test_abs(int x) {
// CHECK-LABEL: @test_labs(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[ABS_I:%.*]] = tail call noundef i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true)
-// CHECK-NEXT: ret i64 [[ABS_I]]
+// CHECK-NEXT: [[TMP0:%.*]] = tail call noundef i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true)
+// CHECK-NEXT: ret i64 [[TMP0]]
//
extern "C" __device__ long test_labs(long x) {
return labs(x);
@@ -249,8 +249,8 @@ extern "C" __device__ long test_labs(long x) {
// CHECK-LABEL: @test_llabs(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[ABS_I:%.*]] = tail call noundef i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true)
-// CHECK-NEXT: ret i64 [[ABS_I]]
+// CHECK-NEXT: [[TMP0:%.*]] = tail call noundef i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true)
+// CHECK-NEXT: ret i64 [[TMP0]]
//
extern "C" __device__ long long test_llabs(long x) {
return llabs(x);
@@ -649,8 +649,8 @@ extern "C" __device__ double test_copysign(double x, double y) {
//
// APPROX-LABEL: @test_cosf(
// APPROX-NEXT: entry:
-// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR16:[0-9]+]]
-// APPROX-NEXT: ret float [[CALL_I_I]]
+// APPROX-NEXT: [[CALL_I1:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR16:[0-9]+]]
+// APPROX-NEXT: ret float [[CALL_I1]]
//
extern "C" __device__ float test_cosf(float x) {
return cosf(x);
@@ -1670,30 +1670,30 @@ extern "C" __device__ double test_j1(double x) {
// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// DEFAULT-NEXT: ]
// DEFAULT: if.then.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
// DEFAULT-NEXT: br label [[_ZL3JNFIF_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR16]]
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
-// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
// DEFAULT: for.body.i:
-// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
-// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]]
-// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]]
-// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
+// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
+// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
+// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// DEFAULT-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
// DEFAULT: _ZL3jnfif.exit:
-// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// DEFAULT-NEXT: ret float [[RETVAL_0_I]]
//
// FINITEONLY-LABEL: @test_jnf(
@@ -1703,30 +1703,30 @@ extern "C" __device__ double test_j1(double x) {
// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// FINITEONLY-NEXT: ]
// FINITEONLY: if.then.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j0_f32(float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j0_f32(float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
// FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j0_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
-// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j0_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
// FINITEONLY: for.body.i:
-// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]]
-// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_025_I]], [[DIV_I]]
-// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_024_I]]
-// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
+// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_0_I3]], [[DIV_I]]
+// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_0_I2]]
+// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// FINITEONLY-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
// FINITEONLY: _ZL3jnfif.exit:
-// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// FINITEONLY-NEXT: ret float [[RETVAL_0_I]]
//
// APPROX-LABEL: @test_jnf(
@@ -1736,30 +1736,30 @@ extern "C" __device__ double test_j1(double x) {
// APPROX-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// APPROX-NEXT: ]
// APPROX: if.then.i:
-// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
+// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
// APPROX-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]]
// APPROX: if.then2.i:
-// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
// APPROX-NEXT: br label [[_ZL3JNFIF_EXIT]]
// APPROX: if.end4.i:
-// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR16]]
-// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
-// APPROX-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// APPROX-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
+// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// APPROX-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
// APPROX: for.body.i:
-// APPROX-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// APPROX-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// APPROX-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// APPROX-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// APPROX-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// APPROX-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// APPROX-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
-// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]]
-// APPROX-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]]
-// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// APPROX-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// APPROX-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
+// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
+// APPROX-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
+// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// APPROX-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// APPROX-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
// APPROX: _ZL3jnfif.exit:
-// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// APPROX-NEXT: ret float [[RETVAL_0_I]]
//
extern "C" __device__ float test_jnf(int x, float y) {
@@ -1773,30 +1773,30 @@ extern "C" __device__ float test_jnf(int x, float y) {
// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// DEFAULT-NEXT: ]
// DEFAULT: if.then.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: br label [[_ZL2JNID_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
// DEFAULT-NEXT: br label [[_ZL2JNID_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR16]]
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
-// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
// DEFAULT: for.body.i:
-// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
-// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]]
-// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]]
-// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP15:![0-9]+]]
+// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
+// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]
+// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// DEFAULT-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP15:![0-9]+]]
// DEFAULT: _ZL2jnid.exit:
-// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// DEFAULT-NEXT: ret double [[RETVAL_0_I]]
//
// FINITEONLY-LABEL: @test_jn(
@@ -1806,30 +1806,30 @@ extern "C" __device__ float test_jnf(int x, float y) {
// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// FINITEONLY-NEXT: ]
// FINITEONLY: if.then.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j0_f64(double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j0_f64(double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
// FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j0_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
-// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j0_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
// FINITEONLY: for.body.i:
-// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]]
-// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_025_I]], [[DIV_I]]
-// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_024_I]]
-// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP15:![0-9]+]]
+// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_0_I3]], [[DIV_I]]
+// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_0_I2]]
+// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// FINITEONLY-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP15:![0-9]+]]
// FINITEONLY: _ZL2jnid.exit:
-// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// FINITEONLY-NEXT: ret double [[RETVAL_0_I]]
//
// APPROX-LABEL: @test_jn(
@@ -1839,30 +1839,30 @@ extern "C" __device__ float test_jnf(int x, float y) {
// APPROX-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// APPROX-NEXT: ]
// APPROX: if.then.i:
-// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
+// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
// APPROX-NEXT: br label [[_ZL2JNID_EXIT:%.*]]
// APPROX: if.then2.i:
-// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
// APPROX-NEXT: br label [[_ZL2JNID_EXIT]]
// APPROX: if.end4.i:
-// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR16]]
-// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
-// APPROX-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// APPROX-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
+// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// APPROX-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
// APPROX: for.body.i:
-// APPROX-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// APPROX-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// APPROX-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// APPROX-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// APPROX-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// APPROX-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// APPROX-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
-// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]]
-// APPROX-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]]
-// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// APPROX-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// APPROX-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP15:![0-9]+]]
+// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
+// APPROX-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]
+// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// APPROX-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// APPROX-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP15:![0-9]+]]
// APPROX: _ZL2jnid.exit:
-// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// APPROX-NEXT: ret double [[RETVAL_0_I]]
//
extern "C" __device__ double test_jn(int x, double y) {
@@ -2364,26 +2364,26 @@ extern "C" __device__ double test_modf(double x, double* y) {
// 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_I30_I_I:%.*]]
+// 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_I14_I_I:%.*]] [
-// CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_I_PREHEADER:%.*]]
-// CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_I_PREHEADER]]
+// 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.i.i.i.preheader:
-// CHECK-NEXT: br label [[WHILE_COND_I_I_I:%.*]]
-// 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]], [[WHILE_COND_I_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[WHILE_COND_I_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA4]]
-// CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP2]], 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: 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_I_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
+// 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
@@ -2391,33 +2391,54 @@ extern "C" __device__ double test_modf(double x, double* y) {
// 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_I_I_I]]
+// 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_I_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 4
+// 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:%.*]] = sext 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:%.*]] = sext 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_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I_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_I_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_I_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP11]]
+// 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:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
-// CHECK-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]], [[CLEANUP_I20_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
-// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[TBAA4]]
-// CHECK-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// 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: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// CHECK-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// 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:%.*]] = shl i64 [[__R_0_I16_I_I]], 3
-// CHECK-NEXT: [[CONV5_I26_I_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
+// CHECK-NEXT: [[CONV5_I26_I_I:%.*]] = sext 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
@@ -2425,30 +2446,9 @@ extern "C" __device__ double test_modf(double x, double* y) {
// 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 [[LOOP7]]
-// 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:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
-// CHECK-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_1_I38_I_I:%.*]], [[CLEANUP_I36_I_I]] ], [ 0, [[ENTRY]] ]
-// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align 1, !tbaa [[TBAA4]]
-// CHECK-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP8]], 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: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// CHECK-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I35_I_I]], label [[IF_THEN_I40_I_I:%.*]], label [[CLEANUP_I36_I_I]]
-// CHECK: if.then.i40.i.i:
-// CHECK-NEXT: [[MUL_I41_I_I:%.*]] = mul i64 [[__R_0_I32_I_I]], 10
-// CHECK-NEXT: [[CONV5_I42_I_I:%.*]] = sext i8 [[TMP8]] to i64
-// CHECK-NEXT: [[ADD_I43_I_I:%.*]] = add i64 [[MUL_I41_I_I]], -48
-// CHECK-NEXT: [[SUB_I44_I_I:%.*]] = add i64 [[ADD_I43_I_I]], [[CONV5_I42_I_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I45_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_I45_I_I]], [[IF_THEN_I40_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]], [[WHILE_BODY_I34_I_I]] ]
-// CHECK-NEXT: [[__R_1_I38_I_I]] = phi i64 [ [[SUB_I44_I_I]], [[IF_THEN_I40_I_I]] ], [ [[__R_0_I32_I_I]], [[WHILE_BODY_I34_I_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I35_I_I]], label [[WHILE_COND_I30_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP10]]
+// 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_I20_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_COND_I14_I_I]] ], [ 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]] ]
+// 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 i32 [[BF_VALUE_I]], 2143289344
@@ -2463,26 +2463,26 @@ extern "C" __device__ float test_nanf(const char *tag) {
// 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_I30_I_I:%.*]]
+// 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_I14_I_I:%.*]] [
-// CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_I_PREHEADER:%.*]]
-// CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_I_PREHEADER]]
+// 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.i.i.i.preheader:
-// CHECK-NEXT: br label [[WHILE_COND_I_I_I:%.*]]
-// 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]], [[WHILE_COND_I_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[WHILE_COND_I_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA4]]
-// CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP2]], 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: 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_I_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
+// 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
@@ -2490,33 +2490,54 @@ extern "C" __device__ float test_nanf(const char *tag) {
// 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_I_I_I]]
+// 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_I_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 4
+// 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:%.*]] = sext 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:%.*]] = sext 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_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I_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_I_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_I_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP11]]
+// 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:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
-// CHECK-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]], [[CLEANUP_I20_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
-// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[TBAA4]]
-// CHECK-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// 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: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// CHECK-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// 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:%.*]] = shl i64 [[__R_0_I16_I_I]], 3
-// CHECK-NEXT: [[CONV5_I26_I_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
+// CHECK-NEXT: [[CONV5_I26_I_I:%.*]] = sext 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
@@ -2524,30 +2545,9 @@ extern "C" __device__ float test_nanf(const char *tag) {
// 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 [[LOOP7]]
-// 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:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
-// CHECK-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_1_I38_I_I:%.*]], [[CLEANUP_I36_I_I]] ], [ 0, [[ENTRY]] ]
-// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align 1, !tbaa [[TBAA4]]
-// CHECK-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP8]], 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: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// CHECK-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I35_I_I]], label [[IF_THEN_I40_I_I:%.*]], label [[CLEANUP_I36_I_I]]
-// CHECK: if.then.i40.i.i:
-// CHECK-NEXT: [[MUL_I41_I_I:%.*]] = mul i64 [[__R_0_I32_I_I]], 10
-// CHECK-NEXT: [[CONV5_I42_I_I:%.*]] = sext i8 [[TMP8]] to i64
-// CHECK-NEXT: [[ADD_I43_I_I:%.*]] = add i64 [[MUL_I41_I_I]], -48
-// CHECK-NEXT: [[SUB_I44_I_I:%.*]] = add i64 [[ADD_I43_I_I]], [[CONV5_I42_I_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I45_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_I45_I_I]], [[IF_THEN_I40_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]], [[WHILE_BODY_I34_I_I]] ]
-// CHECK-NEXT: [[__R_1_I38_I_I]] = phi i64 [ [[SUB_I44_I_I]], [[IF_THEN_I40_I_I]] ], [ [[__R_0_I32_I_I]], [[WHILE_BODY_I34_I_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I35_I_I]], label [[WHILE_COND_I30_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP10]]
+// 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_I20_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_COND_I14_I_I]] ], [ 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]] ]
+// 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 i64 [[BF_VALUE_I]], 9221120237041090560
// CHECK-NEXT: [[TMP10:%.*]] = bitcast i64 [[BF_SET9_I]] to double
@@ -2819,62 +2819,62 @@ extern "C" __device__ double test_normcdfinv(double x) {
// DEFAULT-LABEL: @test_normf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// DEFAULT: while.body.i:
-// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA16]]
+// DEFAULT-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA16]]
// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
-// DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]]
-// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1
// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
// DEFAULT: _ZL5normfiPKf.exit:
-// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[__R_0_LCSSA_I]])
+// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[__R_0_I_LCSSA]])
// DEFAULT-NEXT: ret float [[TMP1]]
//
// FINITEONLY-LABEL: @test_normf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// FINITEONLY: while.body.i:
-// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA16]]
+// FINITEONLY-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA16]]
// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]]
-// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_08_I]], [[MUL_I]]
-// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_0_I4]], [[MUL_I]]
+// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1
// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
// FINITEONLY: _ZL5normfiPKf.exit:
-// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract noundef float @llvm.sqrt.f32(float [[__R_0_LCSSA_I]])
+// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract noundef float @llvm.sqrt.f32(float [[__R_0_I_LCSSA]])
// FINITEONLY-NEXT: ret float [[TMP1]]
//
// APPROX-LABEL: @test_normf(
// APPROX-NEXT: entry:
-// APPROX-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// APPROX-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// APPROX-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// APPROX-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// APPROX: while.body.i:
-// APPROX-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// APPROX-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// APPROX-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// APPROX-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA16]]
+// APPROX-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// APPROX-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// APPROX-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// APPROX-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA16]]
// APPROX-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
-// APPROX-NEXT: [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]]
-// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// APPROX-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1
// APPROX-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// APPROX-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
// APPROX: _ZL5normfiPKf.exit:
-// APPROX-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[__R_0_LCSSA_I]])
+// APPROX-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[__R_0_I_LCSSA]])
// APPROX-NEXT: ret float [[TMP1]]
//
extern "C" __device__ float test_normf(int x, const float *y) {
@@ -2883,62 +2883,62 @@ extern "C" __device__ float test_normf(int x, const float *y) {
// DEFAULT-LABEL: @test_norm(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// DEFAULT: while.body.i:
-// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA18]]
+// DEFAULT-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA18]]
// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
-// DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]]
-// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1
// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
// DEFAULT: _ZL4normiPKd.exit:
-// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[__R_0_LCSSA_I]])
+// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[__R_0_I_LCSSA]])
// DEFAULT-NEXT: ret double [[TMP1]]
//
// FINITEONLY-LABEL: @test_norm(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// FINITEONLY: while.body.i:
-// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA18]]
+// FINITEONLY-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA18]]
// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]]
-// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_08_I]], [[MUL_I]]
-// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_0_I4]], [[MUL_I]]
+// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1
// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
// FINITEONLY: _ZL4normiPKd.exit:
-// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract noundef double @llvm.sqrt.f64(double [[__R_0_LCSSA_I]])
+// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract noundef double @llvm.sqrt.f64(double [[__R_0_I_LCSSA]])
// FINITEONLY-NEXT: ret double [[TMP1]]
//
// APPROX-LABEL: @test_norm(
// APPROX-NEXT: entry:
-// APPROX-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// APPROX-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// APPROX-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// APPROX-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// APPROX: while.body.i:
-// APPROX-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// APPROX-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// APPROX-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// APPROX-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA18]]
+// APPROX-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// APPROX-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// APPROX-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// APPROX-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA18]]
// APPROX-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
-// APPROX-NEXT: [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]]
-// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// APPROX-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1
// APPROX-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// APPROX-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
// APPROX: _ZL4normiPKd.exit:
-// APPROX-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[__R_0_LCSSA_I]])
+// APPROX-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[__R_0_I_LCSSA]])
// APPROX-NEXT: ret double [[TMP1]]
//
extern "C" __device__ double test_norm(int x, const double *y) {
@@ -3243,62 +3243,62 @@ extern "C" __device__ double test_rint(double x) {
// DEFAULT-LABEL: @test_rnormf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// DEFAULT: while.body.i:
-// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA16]]
+// DEFAULT-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA16]]
// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
-// DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]]
-// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1
// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
// DEFAULT: _ZL6rnormfiPKf.exit:
-// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rsqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR15]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rnormf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// FINITEONLY: while.body.i:
-// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA16]]
+// FINITEONLY-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA16]]
// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]]
-// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_08_I]], [[MUL_I]]
-// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_0_I4]], [[MUL_I]]
+// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1
// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
// FINITEONLY: _ZL6rnormfiPKf.exit:
-// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_rsqrt_f32(float noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_rsqrt_f32(float noundef nofpclass(nan inf) [[__R_0_I_LCSSA]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
// APPROX-LABEL: @test_rnormf(
// APPROX-NEXT: entry:
-// APPROX-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// APPROX-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// APPROX-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// APPROX-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// APPROX: while.body.i:
-// APPROX-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// APPROX-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// APPROX-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// APPROX-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA16]]
+// APPROX-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// APPROX-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// APPROX-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// APPROX-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA16]]
// APPROX-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
-// APPROX-NEXT: [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]]
-// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// APPROX-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1
// APPROX-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// APPROX-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
// APPROX: _ZL6rnormfiPKf.exit:
-// APPROX-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// APPROX-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rsqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR15]]
// APPROX-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_rnormf(int x, const float* y) {
@@ -3307,62 +3307,62 @@ extern "C" __device__ float test_rnormf(int x, const float* y) {
// DEFAULT-LABEL: @test_rnorm(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// DEFAULT: while.body.i:
-// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA18]]
+// DEFAULT-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA18]]
// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
-// DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]]
-// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1
// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
// DEFAULT: _ZL5rnormiPKd.exit:
-// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rsqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR15]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rnorm(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// FINITEONLY: while.body.i:
-// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA18]]
+// FINITEONLY-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA18]]
// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]]
-// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_08_I]], [[MUL_I]]
-// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_0_I4]], [[MUL_I]]
+// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1
// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
// FINITEONLY: _ZL5rnormiPKd.exit:
-// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_rsqrt_f64(double noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_rsqrt_f64(double noundef nofpclass(nan inf) [[__R_0_I_LCSSA]]) #[[ATTR15]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
// APPROX-LABEL: @test_rnorm(
// APPROX-NEXT: entry:
-// APPROX-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// APPROX-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// APPROX-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// APPROX-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// APPROX: while.body.i:
-// APPROX-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// APPROX-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// APPROX-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// APPROX-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA18]]
+// APPROX-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// APPROX-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// APPROX-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// APPROX-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA18]]
// APPROX-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
-// APPROX-NEXT: [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]]
-// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// APPROX-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1
// APPROX-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// APPROX-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
// APPROX: _ZL5rnormiPKd.exit:
-// APPROX-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// APPROX-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rsqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR15]]
// APPROX-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_rnorm(int x, const double* y) {
@@ -3838,8 +3838,8 @@ extern "C" __device__ void test_sincospi(double x, double *y, double *z) {
//
// APPROX-LABEL: @test_sinf(
// APPROX-NEXT: entry:
-// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
-// APPROX-NEXT: ret float [[CALL_I_I]]
+// APPROX-NEXT: [[CALL_I1:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT: ret float [[CALL_I1]]
//
extern "C" __device__ float test_sinf(float x) {
return sinf(x);
@@ -4175,30 +4175,30 @@ extern "C" __device__ double test_y1(double x) {
// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// DEFAULT-NEXT: ]
// DEFAULT: if.then.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
// DEFAULT-NEXT: br label [[_ZL3YNFIF_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR16]]
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
-// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
// DEFAULT: for.body.i:
-// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
-// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]]
-// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]]
-// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
+// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
+// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
+// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// DEFAULT-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
// DEFAULT: _ZL3ynfif.exit:
-// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// DEFAULT-NEXT: ret float [[RETVAL_0_I]]
//
// FINITEONLY-LABEL: @test_ynf(
@@ -4208,30 +4208,30 @@ extern "C" __device__ double test_y1(double x) {
// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// FINITEONLY-NEXT: ]
// FINITEONLY: if.then.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y0_f32(float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y0_f32(float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
// FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y0_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
-// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y0_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
// FINITEONLY: for.body.i:
-// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]]
-// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_025_I]], [[DIV_I]]
-// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_024_I]]
-// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
+// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_0_I3]], [[DIV_I]]
+// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_0_I2]]
+// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// FINITEONLY-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
// FINITEONLY: _ZL3ynfif.exit:
-// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// FINITEONLY-NEXT: ret float [[RETVAL_0_I]]
//
// APPROX-LABEL: @test_ynf(
@@ -4241,30 +4241,30 @@ extern "C" __device__ double test_y1(double x) {
// APPROX-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// APPROX-NEXT: ]
// APPROX: if.then.i:
-// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
+// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
// APPROX-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]]
// APPROX: if.then2.i:
-// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
// APPROX-NEXT: br label [[_ZL3YNFIF_EXIT]]
// APPROX: if.end4.i:
-// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR16]]
-// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
-// APPROX-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// APPROX-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
+// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// APPROX-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
// APPROX: for.body.i:
-// APPROX-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// APPROX-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// APPROX-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// APPROX-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// APPROX-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// APPROX-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// APPROX-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
-// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]]
-// APPROX-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]]
-// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// APPROX-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// APPROX-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
+// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
+// APPROX-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
+// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// APPROX-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// APPROX-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
// APPROX: _ZL3ynfif.exit:
-// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// APPROX-NEXT: ret float [[RETVAL_0_I]]
//
extern "C" __device__ float test_ynf(int x, float y) {
@@ -4278,30 +4278,30 @@ extern "C" __device__ float test_ynf(int x, float y) {
// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// DEFAULT-NEXT: ]
// DEFAULT: if.then.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
// DEFAULT-NEXT: br label [[_ZL2YNID_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
// DEFAULT-NEXT: br label [[_ZL2YNID_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR16]]
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
-// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
// DEFAULT: for.body.i:
-// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
-// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]]
-// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]]
-// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]]
+// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
+// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]
+// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// DEFAULT-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]]
// DEFAULT: _ZL2ynid.exit:
-// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// DEFAULT-NEXT: ret double [[RETVAL_0_I]]
//
// FINITEONLY-LABEL: @test_yn(
@@ -4311,30 +4311,30 @@ extern "C" __device__ float test_ynf(int x, float y) {
// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// FINITEONLY-NEXT: ]
// FINITEONLY: if.then.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y0_f64(double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y0_f64(double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]]
// FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
// FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y0_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
-// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y0_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
// FINITEONLY: for.body.i:
-// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]]
-// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_025_I]], [[DIV_I]]
-// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_024_I]]
-// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]]
+// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_0_I3]], [[DIV_I]]
+// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_0_I2]]
+// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// FINITEONLY-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]]
// FINITEONLY: _ZL2ynid.exit:
-// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// FINITEONLY-NEXT: ret double [[RETVAL_0_I]]
//
// APPROX-LABEL: @test_yn(
@@ -4344,30 +4344,30 @@ extern "C" __device__ float test_ynf(int x, float y) {
// APPROX-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// APPROX-NEXT: ]
// APPROX: if.then.i:
-// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
+// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
// APPROX-NEXT: br label [[_ZL2YNID_EXIT:%.*]]
// APPROX: if.then2.i:
-// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
// APPROX-NEXT: br label [[_ZL2YNID_EXIT]]
// APPROX: if.end4.i:
-// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR16]]
-// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
-// APPROX-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// APPROX-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
+// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// APPROX-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
// APPROX: for.body.i:
-// APPROX-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// APPROX-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// APPROX-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// APPROX-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// APPROX-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// APPROX-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// APPROX-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
-// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]]
-// APPROX-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]]
-// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// APPROX-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// APPROX-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]]
+// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
+// APPROX-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]
+// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// APPROX-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// APPROX-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]]
// APPROX: _ZL2ynid.exit:
-// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// APPROX-NEXT: ret double [[RETVAL_0_I]]
//
extern "C" __device__ double test_yn(int x, double y) {
@@ -4742,26 +4742,26 @@ extern "C" __device__ float test___sinf(float x) {
// DEFAULT-LABEL: @test___tanf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
-// DEFAULT-NEXT: [[CALL_I3_I:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]]
-// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.rcp.f32(float [[CALL_I3_I]])
-// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[CALL_I_I]], [[TMP0]]
+// DEFAULT-NEXT: [[CALL_I3_I:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]]
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.rcp.f32(float [[CALL_I_I]])
+// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[CALL_I3_I]], [[TMP0]]
// DEFAULT-NEXT: ret float [[MUL_I]]
//
// FINITEONLY-LABEL: @test___tanf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_native_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
-// FINITEONLY-NEXT: [[CALL_I3_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_native_cos_f32(float noundef nofpclass(nan inf) [[X]]) #[[ATTR16]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.amdgcn.rcp.f32(float [[CALL_I3_I]])
-// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[CALL_I_I]], [[TMP0]]
+// FINITEONLY-NEXT: [[CALL_I3_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_native_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_native_cos_f32(float noundef nofpclass(nan inf) [[X]]) #[[ATTR16]]
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.amdgcn.rcp.f32(float [[CALL_I_I]])
+// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[CALL_I3_I]], [[TMP0]]
// FINITEONLY-NEXT: ret float [[MUL_I]]
//
// APPROX-LABEL: @test___tanf(
// APPROX-NEXT: entry:
-// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
-// APPROX-NEXT: [[CALL_I3_I:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]]
-// APPROX-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.rcp.f32(float [[CALL_I3_I]])
-// APPROX-NEXT: [[MUL_I:%.*]] = fmul contract float [[CALL_I_I]], [[TMP0]]
+// APPROX-NEXT: [[CALL_I3_I:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]]
+// APPROX-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.rcp.f32(float [[CALL_I_I]])
+// APPROX-NEXT: [[MUL_I:%.*]] = fmul contract float [[CALL_I3_I]], [[TMP0]]
// APPROX-NEXT: ret float [[MUL_I]]
//
extern "C" __device__ float test___tanf(float x) {
diff --git a/clang/test/OpenMP/bug57757.cpp b/clang/test/OpenMP/bug57757.cpp
index 7894796ac46284c..7acfe134ddd0baf 100644
--- a/clang/test/OpenMP/bug57757.cpp
+++ b/clang/test/OpenMP/bug57757.cpp
@@ -32,24 +32,23 @@ void foo() {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [[STRUCT_KMP_TASK_T:%.*]], ptr [[TMP1]], i64 0, i32 2
// CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META13:![0-9]+]])
-// CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META16:![0-9]+]])
-// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4, !tbaa [[TBAA18:![0-9]+]], !alias.scope !13, !noalias !16
+// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4, !tbaa [[TBAA16:![0-9]+]], !alias.scope !13, !noalias !17
// CHECK-NEXT: switch i32 [[TMP3]], label [[DOTOMP_OUTLINED__EXIT:%.*]] [
// CHECK-NEXT: i32 0, label [[DOTUNTIED_JMP__I:%.*]]
// CHECK-NEXT: i32 1, label [[DOTUNTIED_NEXT__I:%.*]]
// CHECK-NEXT: ]
// CHECK: .untied.jmp..i:
-// CHECK-NEXT: store i32 1, ptr [[TMP2]], align 4, !tbaa [[TBAA18]], !alias.scope !13, !noalias !16
-// CHECK-NEXT: [[TMP4:%.*]] = tail call i32 @__kmpc_omp_task(ptr nonnull @[[GLOB1]], i32 [[TMP0]], ptr [[TMP1]]), !noalias !19
+// CHECK-NEXT: store i32 1, ptr [[TMP2]], align 4, !tbaa [[TBAA16]], !alias.scope !13, !noalias !17
+// CHECK-NEXT: [[TMP4:%.*]] = tail call i32 @__kmpc_omp_task(ptr nonnull @[[GLOB1]], i32 [[TMP0]], ptr [[TMP1]]), !noalias !13
// CHECK-NEXT: br label [[DOTOMP_OUTLINED__EXIT]]
// CHECK: .untied.next..i:
// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [[STRUCT_KMP_TASK_T_WITH_PRIVATES:%.*]], ptr [[TMP1]], i64 0, i32 1
// CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [[STRUCT_KMP_TASK_T_WITH_PRIVATES]], ptr [[TMP1]], i64 0, i32 1, i32 2
// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds [[STRUCT_KMP_TASK_T_WITH_PRIVATES]], ptr [[TMP1]], i64 0, i32 1, i32 1
-// CHECK-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP5]], align 8, !tbaa [[TBAA20:![0-9]+]], !alias.scope !16, !noalias !13
-// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP7]], align 4, !tbaa [[TBAA18]], !alias.scope !16, !noalias !13
-// CHECK-NEXT: [[TMP10:%.*]] = load float, ptr [[TMP6]], align 4, !tbaa [[TBAA21:![0-9]+]], !alias.scope !16, !noalias !13
-// CHECK-NEXT: tail call void [[TMP8]](i32 noundef [[TMP9]], float noundef [[TMP10]]) #[[ATTR2:[0-9]+]], !noalias !19
+// CHECK-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP5]], align 8, !tbaa [[TBAA19:![0-9]+]], !noalias !13
+// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP7]], align 4, !tbaa [[TBAA16]], !noalias !13
+// CHECK-NEXT: [[TMP10:%.*]] = load float, ptr [[TMP6]], align 4, !tbaa [[TBAA20:![0-9]+]], !noalias !13
+// CHECK-NEXT: tail call void [[TMP8]](i32 noundef [[TMP9]], float noundef [[TMP10]]) #[[ATTR2:[0-9]+]], !noalias !13
// CHECK-NEXT: br label [[DOTOMP_OUTLINED__EXIT]]
// CHECK: .omp_outlined..exit:
// CHECK-NEXT: ret i32 0
diff --git a/llvm/lib/Passes/PassBuilderPipelines.cpp b/llvm/lib/Passes/PassBuilderPipelines.cpp
index 600f8d43caaf216..baea2913338cda7 100644
--- a/llvm/lib/Passes/PassBuilderPipelines.cpp
+++ b/llvm/lib/Passes/PassBuilderPipelines.cpp
@@ -164,7 +164,7 @@ static cl::opt<bool> EnableModuleInliner("enable-module-inliner",
cl::desc("Enable module inliner"));
static cl::opt<bool> PerformMandatoryInliningsFirst(
- "mandatory-inlining-first", cl::init(true), cl::Hidden,
+ "mandatory-inlining-first", cl::init(false), cl::Hidden,
cl::desc("Perform mandatory inlinings module-wide, before performing "
"inlining"));
@@ -1127,6 +1127,8 @@ PassBuilder::buildModuleSimplificationPipeline(OptimizationLevel Level,
if (EnableSyntheticCounts && !PGOOpt)
MPM.addPass(SyntheticCountsPropagation());
+ MPM.addPass(AlwaysInlinerPass(/*InsertLifetimeIntrinsics=*/true));
+
if (EnableModuleInliner)
MPM.addPass(buildModuleInlinerPipeline(Level, Phase));
else
diff --git a/llvm/test/Other/new-pm-defaults.ll b/llvm/test/Other/new-pm-defaults.ll
index 098db2b959a29ec..ecdb5a5e010d922 100644
--- a/llvm/test/Other/new-pm-defaults.ll
+++ b/llvm/test/Other/new-pm-defaults.ll
@@ -121,6 +121,8 @@
; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy
; CHECK-EP-PEEPHOLE-NEXT: Running pass: NoOpFunctionPass
; CHECK-O-NEXT: Running pass: SimplifyCFGPass
+; CHECK-O-NEXT: Running pass: AlwaysInlinerPass
+; CHECK-O-NEXT: Running analysis: ProfileSummaryAnalysis
; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass
; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis
; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA
@@ -129,14 +131,12 @@
; CHECK-O-NEXT: Running pass: InvalidateAnalysisPass<{{.*}}AAManager
; CHECK-O-NEXT: Invalidating analysis: AAManager
; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}ProfileSummaryAnalysis
-; CHECK-O-NEXT: Running analysis: ProfileSummaryAnalysis
; CHECK-O-NEXT: Running analysis: InnerAnalysisManagerProxy
; CHECK-O-NEXT: Running analysis: LazyCallGraphAnalysis
; CHECK-O-NEXT: Running analysis: FunctionAnalysisManagerCGSCCProxy
; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy<{{.*}}LazyCallGraph::SCC{{.*}}>
; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass
; CHECK-O-NEXT: Running pass: InlinerPass
-; CHECK-O-NEXT: Running pass: InlinerPass
; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass
; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass on (foo)
diff --git a/llvm/test/Other/new-pm-print-pipeline.ll b/llvm/test/Other/new-pm-print-pipeline.ll
index 6214d6abcef3a36..c9b06ac7e6e8e72 100644
--- a/llvm/test/Other/new-pm-print-pipeline.ll
+++ b/llvm/test/Other/new-pm-print-pipeline.ll
@@ -62,7 +62,7 @@
; CHECK-20: cgscc(inline<only-mandatory>,inline),cgscc(inline)
; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='scc-oz-module-inliner' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-21
-; CHECK-21: require<globals-aa>,function(invalidate<aa>),require<profile-summary>,cgscc(devirt<4>(inline<only-mandatory>,inline,{{.*}},instcombine{{.*}}))
+; CHECK-21: require<globals-aa>,function(invalidate<aa>),require<profile-summary>,cgscc(devirt<4>(inline,{{.*}},instcombine{{.*}}))
; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='cgscc(function<eager-inv>(no-op-function)),function<eager-inv>(no-op-function)' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-22
; CHECK-22: cgscc(function<eager-inv>(no-op-function)),function<eager-inv>(no-op-function)
diff --git a/llvm/test/Other/new-pm-thinlto-postlink-defaults.ll b/llvm/test/Other/new-pm-thinlto-postlink-defaults.ll
index 1103c8f984fa296..064362eabbf8395 100644
--- a/llvm/test/Other/new-pm-thinlto-postlink-defaults.ll
+++ b/llvm/test/Other/new-pm-thinlto-postlink-defaults.ll
@@ -61,6 +61,8 @@
; CHECK-O-NEXT: Running analysis: TypeBasedAA
; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy
; CHECK-O-NEXT: Running pass: SimplifyCFGPass
+; CHECK-O-NEXT: Running pass: AlwaysInlinerPass
+; CHECK-PRELINK-O-NEXT: Running analysis: ProfileSummaryAnalysis
; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass
; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis
; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA
@@ -75,7 +77,6 @@
; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy
; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass
; CHECK-O-NEXT: Running pass: InlinerPass
-; CHECK-O-NEXT: Running pass: InlinerPass
; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass
; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass on (foo)
diff --git a/llvm/test/Other/new-pm-thinlto-postlink-pgo-defaults.ll b/llvm/test/Other/new-pm-thinlto-postlink-pgo-defaults.ll
index 05d165ae399f903..19a44867e434ace 100644
--- a/llvm/test/Other/new-pm-thinlto-postlink-pgo-defaults.ll
+++ b/llvm/test/Other/new-pm-thinlto-postlink-pgo-defaults.ll
@@ -50,6 +50,7 @@
; CHECK-O-NEXT: Running analysis: LoopAnalysis on foo
; CHECK-O-NEXT: Running analysis: PostDominatorTreeAnalysis on foo
; CHECK-O-NEXT: Running pass: SimplifyCFGPass
+; CHECK-O-NEXT: Running pass: AlwaysInlinerPass
; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass
; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis
; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA
@@ -64,7 +65,6 @@
; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy<{{.*}}LazyCallGraph::SCC{{.*}}>
; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass
; CHECK-O-NEXT: Running pass: InlinerPass
-; CHECK-O-NEXT: Running pass: InlinerPass
; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass
; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass
diff --git a/llvm/test/Other/new-pm-thinlto-postlink-samplepgo-defaults.ll b/llvm/test/Other/new-pm-thinlto-postlink-samplepgo-defaults.ll
index 2393ba9dbde4b64..ac80a31d8fd4bc0 100644
--- a/llvm/test/Other/new-pm-thinlto-postlink-samplepgo-defaults.ll
+++ b/llvm/test/Other/new-pm-thinlto-postlink-samplepgo-defaults.ll
@@ -58,6 +58,8 @@
; CHECK-O-NEXT: Running analysis: LoopAnalysis on foo
; CHECK-O-NEXT: Running analysis: PostDominatorTreeAnalysis on foo
; CHECK-O-NEXT: Running pass: SimplifyCFGPass on foo
+
+; CHECK-O-NEXT: Running pass: AlwaysInlinerPass
; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass
; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis
; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA
@@ -71,7 +73,6 @@
; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy
; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass
; CHECK-O-NEXT: Running pass: InlinerPass
-; CHECK-O-NEXT: Running pass: InlinerPass
; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass
; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass
diff --git a/llvm/test/Other/new-pm-thinlto-prelink-defaults.ll b/llvm/test/Other/new-pm-thinlto-prelink-defaults.ll
index 722d586c86907e7..6486639e07b49c0 100644
--- a/llvm/test/Other/new-pm-thinlto-prelink-defaults.ll
+++ b/llvm/test/Other/new-pm-thinlto-prelink-defaults.ll
@@ -92,6 +92,8 @@
; CHECK-O-NEXT: Running analysis: TypeBasedAA
; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy
; CHECK-O-NEXT: Running pass: SimplifyCFGPass
+; CHECK-O-NEXT: Running pass: AlwaysInlinerPass
+; CHECK-O-NEXT: Running analysis: ProfileSummaryAnalysis
; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass
; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis
; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA
@@ -100,14 +102,12 @@
; CHECK-O-NEXT: Running pass: InvalidateAnalysisPass<{{.*}}AAManager
; CHECK-O-NEXT: Invalidating analysis: AAManager
; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}ProfileSummaryAnalysis
-; CHECK-O-NEXT: Running analysis: ProfileSummaryAnalysis
; CHECK-O-NEXT: Running analysis: InnerAnalysisManagerProxy
; CHECK-O-NEXT: Running analysis: LazyCallGraphAnalysis
; CHECK-O-NEXT: Running analysis: FunctionAnalysisManagerCGSCCProxy
; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy
; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass
; CHECK-O-NEXT: Running pass: InlinerPass
-; CHECK-O-NEXT: Running pass: InlinerPass
; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass
; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass on (foo)
diff --git a/llvm/test/Other/new-pm-thinlto-prelink-pgo-defaults.ll b/llvm/test/Other/new-pm-thinlto-prelink-pgo-defaults.ll
index 7e4c82666f0cd30..09f9f0f48baddb2 100644
--- a/llvm/test/Other/new-pm-thinlto-prelink-pgo-defaults.ll
+++ b/llvm/test/Other/new-pm-thinlto-prelink-pgo-defaults.ll
@@ -83,6 +83,7 @@
; CHECK-O-NEXT: Running pass: PGOIndirectCallPromotion on
; CHECK-O-NEXT: Running analysis: InnerAnalysisManagerProxy
; CHECK-O-NEXT: Running analysis: OptimizationRemarkEmitterAnalysis on foo
+; CHECK-O-NEXT: Running pass: AlwaysInlinerPass
; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass
; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis
; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA
@@ -97,7 +98,6 @@
; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy<{{.*}}LazyCallGraph::SCC{{.*}}>
; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass
; CHECK-O-NEXT: Running pass: InlinerPass
-; CHECK-O-NEXT: Running pass: InlinerPass
; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass
; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass
diff --git a/llvm/test/Other/new-pm-thinlto-prelink-samplepgo-defaults.ll b/llvm/test/Other/new-pm-thinlto-prelink-samplepgo-defaults.ll
index c609fbe31de2151..47bdbfd2d357d47 100644
--- a/llvm/test/Other/new-pm-thinlto-prelink-samplepgo-defaults.ll
+++ b/llvm/test/Other/new-pm-thinlto-prelink-samplepgo-defaults.ll
@@ -63,6 +63,7 @@
; CHECK-O-NEXT: Running analysis: LoopAnalysis on foo
; CHECK-O-NEXT: Running analysis: PostDominatorTreeAnalysis on foo
; CHECK-O-NEXT: Running pass: SimplifyCFGPass on foo
+; CHECK-O-NEXT: Running pass: AlwaysInlinerPass
; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass
; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis
; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA
@@ -76,7 +77,6 @@
; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy<{{.*}}LazyCallGraph::SCC{{.*}}>
; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass
; CHECK-O-NEXT: Running pass: InlinerPass
-; CHECK-O-NEXT: Running pass: InlinerPass
; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass
; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass
diff --git a/llvm/test/Transforms/Inline/always-inline-phase-ordering.ll b/llvm/test/Transforms/Inline/always-inline-phase-ordering.ll
new file mode 100644
index 000000000000000..e69ca483449007c
--- /dev/null
+++ b/llvm/test/Transforms/Inline/always-inline-phase-ordering.ll
@@ -0,0 +1,164 @@
+; RUN: opt --Os -pass-remarks=inline -S < %s 2>&1 | FileCheck %s
+target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"
+target triple = "arm64e-apple-macosx13"
+
+; CHECK: remark: <unknown>:0:0: 'wibble' inlined into 'bar.8' with (cost=always): always inline attribute
+; CHECK: remark: <unknown>:0:0: 'wibble' inlined into 'pluto' with (cost=always): always inline attribute
+; CHECK: remark: <unknown>:0:0: 'snork' inlined into 'blam' with (cost=always): always inline attribute
+; CHECK: remark: <unknown>:0:0: 'wobble' inlined into 'blam' with (cost=always): always inline attribute
+; CHECK: remark: <unknown>:0:0: 'wobble' inlined into 'snork' with (cost=always): always inline attribute
+; CHECK: remark: <unknown>:0:0: 'spam' inlined into 'blam' with (cost=65, threshold=75)
+; CHECK: remark: <unknown>:0:0: 'wibble.1' inlined into 'widget' with (cost=30, threshold=75)
+; CHECK: remark: <unknown>:0:0: 'widget' inlined into 'bar.8' with (cost=30, threshold=75)
+; CHECK: remark: <unknown>:0:0: 'barney' inlined into 'wombat' with (cost=30, threshold=75)
+
+define linkonce_odr void @wombat(ptr %arg) #0 {
+bb:
+ call void @barney()
+ ret void
+}
+
+define i1 @foo() {
+bb:
+ call void @wombat(ptr null)
+ unreachable
+}
+
+define linkonce_odr void @pluto() #1 !prof !38 {
+bb:
+ call void @wibble()
+ ret void
+}
+
+; Function Attrs: alwaysinline
+define linkonce_odr void @wibble() #2 {
+bb:
+ call void @widget()
+ ret void
+}
+
+define linkonce_odr void @widget() {
+bb:
+ call void @wibble.1()
+ ret void
+}
+
+define linkonce_odr void @wibble.1() {
+bb:
+ %0 = call i32 @foo.2()
+ call void @blam()
+ ret void
+}
+
+declare i32 @foo.2()
+
+define linkonce_odr void @blam() {
+bb:
+ %tmp = call i32 @snork()
+ %tmpv1 = call ptr @wombat.3()
+ call void @eggs()
+ %tmpv2 = call ptr @wombat.3()
+ ret void
+}
+
+; Function Attrs: alwaysinline
+define linkonce_odr i32 @snork() #2 {
+bb:
+ %tmpv1 = call i32 @spam()
+ %tmpv2 = call i32 @wobble()
+ call void @widget.4(i32 %tmpv2)
+ ret i32 0
+}
+
+declare void @eggs()
+
+declare ptr @wombat.3()
+
+define linkonce_odr i32 @spam() {
+bb:
+ %tmpv1 = call i32 @wombat.6()
+ %tmpv2 = call i64 @wobble.5(i8 0)
+ %tmpv3 = call i64 @bar()
+ ret i32 0
+}
+
+; Function Attrs: alwaysinline
+define linkonce_odr i32 @wobble() #2 {
+bb:
+ %tmpv = call i64 @wobble.5(i8 0)
+ %tmpv1 = call i64 @eggs.7()
+ %tmpv2 = call i64 @wobble.5(i8 0)
+ %tmpv3 = call i64 @eggs.7()
+ %tmpv4 = lshr i64 %tmpv1, 1
+ %tmpv5 = trunc i64 %tmpv4 to i32
+ %tmpv6 = xor i32 %tmpv5, 23
+ ret i32 %tmpv6
+}
+
+declare void @widget.4(i32)
+
+declare i64 @bar()
+
+declare i64 @wobble.5(i8)
+
+declare i32 @wombat.6()
+
+declare i64 @eggs.7()
+
+define linkonce_odr void @barney() {
+bb:
+ call void @bar.8()
+ call void @pluto()
+ unreachable
+}
+
+define linkonce_odr void @bar.8() {
+bb:
+ call void @wibble()
+ ret void
+}
+
+attributes #0 = { "frame-pointer"="non-leaf" }
+attributes #1 = { "target-cpu"="apple-m1" }
+attributes #2 = { alwaysinline }
+
+!llvm.module.flags = !{!0, !1, !30, !31, !32, !36, !37}
+
+!0 = !{i32 2, !"SDK Version", [2 x i32] [i32 13, i32 3]}
+!1 = !{i32 1, !"ProfileSummary", !2}
+!2 = !{!3, !4, !5, !6, !7, !8, !9, !10, !11, !12}
+!3 = !{!"ProfileFormat", !"InstrProf"}
+!4 = !{!"TotalCount", i64 864540306756}
+!5 = !{!"MaxCount", i64 6596759955}
+!6 = !{!"MaxInternalCount", i64 2828618424}
+!7 = !{!"MaxFunctionCount", i64 6596759955}
+!8 = !{!"NumCounts", i64 268920}
+!9 = !{!"NumFunctions", i64 106162}
+!10 = !{!"IsPartialProfile", i64 0}
+!11 = !{!"PartialProfileRatio", double 0.000000e+00}
+!12 = !{!"DetailedSummary", !13}
+!13 = !{!14, !15, !16, !17, !18, !19, !20, !21, !22, !23, !24, !25, !26, !27, !28, !29}
+!14 = !{i32 10000, i64 5109654023, i32 2}
+!15 = !{i32 100000, i64 2480859832, i32 25}
+!16 = !{i32 200000, i64 1566552109, i32 70}
+!17 = !{i32 300000, i64 973667919, i32 140}
+!18 = !{i32 400000, i64 552159773, i32 263}
+!19 = !{i32 500000, i64 353879860, i32 463}
+!20 = !{i32 600000, i64 187122455, i32 799}
+!21 = !{i32 700000, i64 105465980, i32 1419}
+!22 = !{i32 800000, i64 49243829, i32 2620}
+!23 = !{i32 900000, i64 15198227, i32 5898}
+!24 = !{i32 950000, i64 5545670, i32 10696}
+!25 = !{i32 990000, i64 804816, i32 25738}
+!26 = !{i32 999000, i64 73999, i32 53382}
+!27 = !{i32 999900, i64 6530, i32 83503}
+!28 = !{i32 999990, i64 899, i32 110416}
+!29 = !{i32 999999, i64 120, i32 130201}
+!30 = !{i32 7, !"Dwarf Version", i32 4}
+!31 = !{i32 2, !"Debug Info Version", i32 3}
+!32 = !{i32 1, !"wchar_size", i32 4}
+!34 = !{!35}
+!35 = !{i32 0, i1 false}
+!36 = !{i32 8, !"PIC Level", i32 2}
+!37 = !{i32 7, !"frame-pointer", i32 1}
+!38 = !{!"function_entry_count", i64 15128150}
More information about the cfe-commits
mailing list