[clang] 86bfeb9 - Revert "Inlining: Run the legacy AlwaysInliner before the regular inliner."
David Green via cfe-commits
cfe-commits at lists.llvm.org
Fri Feb 10 07:01:56 PST 2023
Author: David Green
Date: 2023-02-10T15:01:49Z
New Revision: 86bfeb906e3a95ae428f3e97d78d3d22a7c839f3
URL: https://github.com/llvm/llvm-project/commit/86bfeb906e3a95ae428f3e97d78d3d22a7c839f3
DIFF: https://github.com/llvm/llvm-project/commit/86bfeb906e3a95ae428f3e97d78d3d22a7c839f3.diff
LOG: Revert "Inlining: Run the legacy AlwaysInliner before the regular inliner."
This seems to cause large regressions in existing code, as much as 75% slower
(4x the time taken). Small always inline functions seem to be used a lot in the
cmsis-dsp library.
I would add a phase ordering test to show the problems, but one already exists!
The llvm/test/Transforms/PhaseOrdering/ARM/arm_mult_q15.ll was just changed by
removing alwaysinline to hide the problems that existed.
This reverts commit cae033dcf227aeecf58fca5af6fc7fde1fd2fb4f.
This reverts commit 8e33c41e72ad42e4c27f8cbc3ad2e02b169637a1.
Added:
Modified:
clang/test/CodeGen/code-coverage.c
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-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-pgo-defaults.ll
llvm/test/Other/new-pm-thinlto-prelink-samplepgo-defaults.ll
llvm/test/Transforms/PhaseOrdering/ARM/arm_mult_q15.ll
Removed:
llvm/test/Transforms/Inline/always-inline-phase-ordering.ll
################################################################################
diff --git a/clang/test/CodeGen/code-coverage.c b/clang/test/CodeGen/code-coverage.c
index ebf3e71672ad7..3748f23cbf674 100644
--- a/clang/test/CodeGen/code-coverage.c
+++ b/clang/test/CodeGen/code-coverage.c
@@ -14,6 +14,7 @@
// RUN: %clang_cc1 -emit-llvm-bc -o /dev/null -fdebug-pass-manager -fprofile-arcs %s 2>&1 | FileCheck --check-prefix=NEWPM %s
// RUN: %clang_cc1 -emit-llvm-bc -o /dev/null -fdebug-pass-manager -fprofile-arcs -O3 %s 2>&1 | FileCheck --check-prefix=NEWPM-O3 %s
+// NEWPM-NOT: Running pass
// NEWPM: Running pass: GCOVProfilerPass
// NEWPM-O3-NOT: Running pass
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 173b43ba41bfe..cb1992ec4cb85 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' with (cost=always): always inline attribute at callsite bar:8:10; (hotness:}}
+ // expected-remark at +1 {{'foo' inlined into 'bar': 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 07eee0abd9e27..edc9731de9bbd 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -130,26 +130,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 [[TBAA3]]
// CHECK-NEXT: [[CMP_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// CHECK-NEXT: br i1 [[CMP_I]], label [[IF_THEN_I:%.*]], label [[WHILE_COND_I17_I:%.*]]
+// CHECK-NEXT: br i1 [[CMP_I]], label [[IF_THEN_I:%.*]], label [[WHILE_COND_I33_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 [[TBAA3]]
-// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I:%.*]] [
-// CHECK-NEXT: i8 120, label [[WHILE_COND_I33_I_PREHEADER:%.*]]
-// CHECK-NEXT: i8 88, label [[WHILE_COND_I33_I_PREHEADER]]
+// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I17_I:%.*]] [
+// CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_PREHEADER:%.*]]
+// CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_PREHEADER]]
// CHECK-NEXT: ]
-// CHECK: while.cond.i33.i.preheader:
-// CHECK-NEXT: br label [[WHILE_COND_I33_I:%.*]]
-// CHECK: while.cond.i33.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I30_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I43_I:%.*]], [[CLEANUP_I44_I:%.*]] ], [ [[INCDEC_PTR_I]], [[WHILE_COND_I33_I_PREHEADER]] ]
-// CHECK-NEXT: [[__R_0_I31_I:%.*]] = phi i64 [ [[__R_2_I_I:%.*]], [[CLEANUP_I44_I]] ], [ 0, [[WHILE_COND_I33_I_PREHEADER]] ]
-// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: [[CMP_NOT_I32_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I32_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT:%.*]], label [[WHILE_BODY_I35_I:%.*]]
-// CHECK: while.body.i35.i:
+// 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 [[TBAA3]]
+// 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-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// CHECK-NEXT: [[OR_COND_I34_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I34_I]], label [[IF_END31_I_I:%.*]], label [[IF_ELSE_I_I:%.*]]
+// 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: if.else.i.i:
// CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
// CHECK-NEXT: [[OR_COND33_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
@@ -157,54 +157,33 @@ 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_I44_I]]
+// CHECK-NEXT: br i1 [[OR_COND34_I_I]], label [[IF_END31_I_I]], label [[CLEANUP_I_I]]
// CHECK: if.end31.i.i:
-// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I35_I]] ], [ -87, [[IF_ELSE_I_I]] ], [ -55, [[IF_ELSE17_I_I]] ]
-// CHECK-NEXT: [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I31_I]], 4
+// 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: [[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_I42_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I]], i64 1
-// CHECK-NEXT: br label [[CLEANUP_I44_I]]
-// CHECK: cleanup.i44.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I43_I]] = phi ptr [ [[INCDEC_PTR_I42_I]], [[IF_END31_I_I]] ], [ [[__TAGP_ADDR_0_I30_I]], [[IF_ELSE17_I_I]] ]
-// CHECK-NEXT: [[__R_2_I_I]] = phi i64 [ [[ADD28_I_I]], [[IF_END31_I_I]] ], [ [[__R_0_I31_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_I33_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP10]]
-// 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 [[TBAA3]]
-// 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_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 [[LOOP6]]
+// 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 [[LOOP10]]
// CHECK: while.cond.i17.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I14_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I:%.*]], [[CLEANUP_I28_I:%.*]] ], [ [[P]], [[ENTRY:%.*]] ]
-// CHECK-NEXT: [[__R_0_I15_I:%.*]] = phi i64 [ [[__R_1_I27_I:%.*]], [[CLEANUP_I28_I]] ], [ 0, [[ENTRY]] ]
-// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: [[CMP_NOT_I16_I:%.*]] = icmp eq i8 [[TMP8]], 0
+// CHECK-NEXT: [[__TAGP_ADDR_0_I14_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I:%.*]], [[CLEANUP_I28_I:%.*]] ], [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ]
+// CHECK-NEXT: [[__R_0_I15_I:%.*]] = phi i64 [ [[__R_1_I27_I:%.*]], [[CLEANUP_I28_I]] ], [ 0, [[IF_THEN_I]] ]
+// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT: [[CMP_NOT_I16_I:%.*]] = icmp eq i8 [[TMP6]], 0
// CHECK-NEXT: br i1 [[CMP_NOT_I16_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I19_I:%.*]]
// CHECK: while.body.i19.i:
-// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// CHECK-NEXT: [[OR_COND_I18_I:%.*]] = icmp ult i8 [[TMP9]], 10
+// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// CHECK-NEXT: [[OR_COND_I18_I:%.*]] = icmp eq i8 [[TMP7]], 48
// CHECK-NEXT: br i1 [[OR_COND_I18_I]], label [[IF_THEN_I25_I:%.*]], label [[CLEANUP_I28_I]]
// CHECK: if.then.i25.i:
-// CHECK-NEXT: [[MUL_I20_I:%.*]] = mul i64 [[__R_0_I15_I]], 10
-// CHECK-NEXT: [[CONV5_I21_I:%.*]] = sext i8 [[TMP8]] to i64
+// CHECK-NEXT: [[MUL_I20_I:%.*]] = shl i64 [[__R_0_I15_I]], 3
+// CHECK-NEXT: [[CONV5_I21_I:%.*]] = sext i8 [[TMP6]] to i64
// CHECK-NEXT: [[ADD_I22_I:%.*]] = add i64 [[MUL_I20_I]], -48
// CHECK-NEXT: [[SUB_I23_I:%.*]] = add i64 [[ADD_I22_I]], [[CONV5_I21_I]]
// CHECK-NEXT: [[INCDEC_PTR_I24_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I14_I]], i64 1
@@ -212,9 +191,30 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
// CHECK: cleanup.i28.i:
// CHECK-NEXT: [[__TAGP_ADDR_1_I26_I]] = phi ptr [ [[INCDEC_PTR_I24_I]], [[IF_THEN_I25_I]] ], [ [[__TAGP_ADDR_0_I14_I]], [[WHILE_BODY_I19_I]] ]
// CHECK-NEXT: [[__R_1_I27_I]] = phi i64 [ [[SUB_I23_I]], [[IF_THEN_I25_I]] ], [ [[__R_0_I15_I]], [[WHILE_BODY_I19_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I18_I]], label [[WHILE_COND_I17_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP9]]
+// CHECK-NEXT: br i1 [[OR_COND_I18_I]], label [[WHILE_COND_I17_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP6]]
+// CHECK: while.cond.i33.i:
+// CHECK-NEXT: [[__TAGP_ADDR_0_I30_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I42_I:%.*]], [[CLEANUP_I44_I:%.*]] ], [ [[P]], [[ENTRY:%.*]] ]
+// CHECK-NEXT: [[__R_0_I31_I:%.*]] = phi i64 [ [[__R_1_I43_I:%.*]], [[CLEANUP_I44_I]] ], [ 0, [[ENTRY]] ]
+// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT: [[CMP_NOT_I32_I:%.*]] = icmp eq i8 [[TMP8]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I32_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I35_I:%.*]]
+// CHECK: while.body.i35.i:
+// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// CHECK-NEXT: [[OR_COND_I34_I:%.*]] = icmp ult i8 [[TMP9]], 10
+// CHECK-NEXT: br i1 [[OR_COND_I34_I]], label [[IF_THEN_I41_I:%.*]], label [[CLEANUP_I44_I]]
+// CHECK: if.then.i41.i:
+// CHECK-NEXT: [[MUL_I36_I:%.*]] = mul i64 [[__R_0_I31_I]], 10
+// CHECK-NEXT: [[CONV5_I37_I:%.*]] = sext i8 [[TMP8]] to i64
+// CHECK-NEXT: [[ADD_I38_I:%.*]] = add i64 [[MUL_I36_I]], -48
+// CHECK-NEXT: [[SUB_I39_I:%.*]] = add i64 [[ADD_I38_I]], [[CONV5_I37_I]]
+// CHECK-NEXT: [[INCDEC_PTR_I40_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I]], i64 1
+// CHECK-NEXT: br label [[CLEANUP_I44_I]]
+// CHECK: cleanup.i44.i:
+// CHECK-NEXT: [[__TAGP_ADDR_1_I42_I]] = phi ptr [ [[INCDEC_PTR_I40_I]], [[IF_THEN_I41_I]] ], [ [[__TAGP_ADDR_0_I30_I]], [[WHILE_BODY_I35_I]] ]
+// CHECK-NEXT: [[__R_1_I43_I]] = phi i64 [ [[SUB_I39_I]], [[IF_THEN_I41_I]] ], [ [[__R_0_I31_I]], [[WHILE_BODY_I35_I]] ]
+// CHECK-NEXT: br i1 [[OR_COND_I34_I]], label [[WHILE_COND_I33_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP9]]
// CHECK: _ZL15__make_mantissaPKc.exit:
-// CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I]] ], [ [[__R_0_I_I]], [[WHILE_COND_I_I]] ], [ 0, [[CLEANUP_I44_I]] ], [ [[__R_0_I31_I]], [[WHILE_COND_I33_I]] ], [ 0, [[CLEANUP_I28_I]] ], [ [[__R_0_I15_I]], [[WHILE_COND_I17_I]] ]
+// CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, [[CLEANUP_I28_I]] ], [ [[__R_0_I15_I]], [[WHILE_COND_I17_I]] ], [ 0, [[CLEANUP_I_I]] ], [ [[__R_0_I_I]], [[WHILE_COND_I_I]] ], [ 0, [[CLEANUP_I44_I]] ], [ [[__R_0_I31_I]], [[WHILE_COND_I33_I]] ]
// CHECK-NEXT: ret i64 [[RETVAL_0_I]]
//
extern "C" __device__ uint64_t test___make_mantissa(const char *p) {
@@ -767,6 +767,7 @@ extern "C" __device__ float test_exp2f(float x) {
return exp2f(x);
}
+//
// DEFAULT-LABEL: @test_exp2(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR14]]
@@ -1284,30 +1285,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_I20_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL3JNFIF_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
// DEFAULT: for.body.i:
-// 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: [[__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: [[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_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 [[LOOP13:![0-9]+]]
+// 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 [[LOOP13:![0-9]+]]
// DEFAULT: _ZL3jnfif.exit:
-// 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: [[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: ret float [[RETVAL_0_I]]
//
// FINITEONLY-LABEL: @test_jnf(
@@ -1317,30 +1318,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_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
// FINITEONLY: for.body.i:
-// 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: [[__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: [[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_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 [[LOOP13:![0-9]+]]
+// 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 [[LOOP13:![0-9]+]]
// FINITEONLY: _ZL3jnfif.exit:
-// 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: [[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: ret float [[RETVAL_0_I]]
//
extern "C" __device__ float test_jnf(int x, float y) {
@@ -1354,30 +1355,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_I20_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL2JNID_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL2JNID_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
// DEFAULT: for.body.i:
-// 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: [[__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: [[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_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 [[LOOP14:![0-9]+]]
+// 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 [[LOOP14:![0-9]+]]
// DEFAULT: _ZL2jnid.exit:
-// 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: [[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: ret double [[RETVAL_0_I]]
//
// FINITEONLY-LABEL: @test_jn(
@@ -1387,30 +1388,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_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
// FINITEONLY: for.body.i:
-// 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: [[__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: [[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_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 [[LOOP14:![0-9]+]]
+// 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 [[LOOP14:![0-9]+]]
// FINITEONLY: _ZL2jnid.exit:
-// 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: [[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: ret double [[RETVAL_0_I]]
//
extern "C" __device__ double test_jn(int x, double y) {
@@ -1765,26 +1766,26 @@ extern "C" __device__ double test_modf(double x, double* y) {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa [[TBAA3]]
// 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_I17_I_I:%.*]]
+// CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I33_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 [[TBAA3]]
-// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [
-// CHECK-NEXT: i8 120, label [[WHILE_COND_I33_I_I_PREHEADER:%.*]]
-// CHECK-NEXT: i8 88, label [[WHILE_COND_I33_I_I_PREHEADER]]
+// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I17_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: ]
-// CHECK: while.cond.i33.i.i.preheader:
-// CHECK-NEXT: br label [[WHILE_COND_I33_I_I:%.*]]
-// CHECK: while.cond.i33.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I30_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I43_I_I:%.*]], [[CLEANUP_I44_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I33_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[__R_0_I31_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I44_I_I]] ], [ 0, [[WHILE_COND_I33_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: [[CMP_NOT_I32_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I32_I_I]], label [[_ZL4NANFPKC_EXIT:%.*]], label [[WHILE_BODY_I35_I_I:%.*]]
-// CHECK: while.body.i35.i.i:
+// 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 [[TBAA3]]
+// 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-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// CHECK-NEXT: [[OR_COND_I34_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
+// 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: if.else.i.i.i:
// CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
// CHECK-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
@@ -1792,54 +1793,33 @@ 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_I44_I_I]]
+// CHECK-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I_I_I]]
// CHECK: if.end31.i.i.i:
-// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I35_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I31_I_I]], 4
+// 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: [[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_I42_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I_I]], i64 1
-// CHECK-NEXT: br label [[CLEANUP_I44_I_I]]
-// CHECK: cleanup.i44.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I43_I_I]] = phi ptr [ [[INCDEC_PTR_I42_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I30_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_I31_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_I33_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP10]]
-// 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 [[TBAA3]]
-// 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_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 [[LOOP6]]
+// 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 [[LOOP10]]
// CHECK: while.cond.i17.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I_I:%.*]], [[CLEANUP_I28_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
-// CHECK-NEXT: [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I27_I_I:%.*]], [[CLEANUP_I28_I_I]] ], [ 0, [[ENTRY]] ]
-// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
+// CHECK-NEXT: [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I_I:%.*]], [[CLEANUP_I28_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
+// CHECK-NEXT: [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I27_I_I:%.*]], [[CLEANUP_I28_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
+// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT: [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
// CHECK-NEXT: br i1 [[CMP_NOT_I16_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I19_I_I:%.*]]
// CHECK: while.body.i19.i.i:
-// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// CHECK-NEXT: [[OR_COND_I18_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
+// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// CHECK-NEXT: [[OR_COND_I18_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
// CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[IF_THEN_I25_I_I:%.*]], label [[CLEANUP_I28_I_I]]
// CHECK: if.then.i25.i.i:
-// CHECK-NEXT: [[MUL_I20_I_I:%.*]] = mul i64 [[__R_0_I15_I_I]], 10
-// CHECK-NEXT: [[CONV5_I21_I_I:%.*]] = sext i8 [[TMP8]] to i64
+// CHECK-NEXT: [[MUL_I20_I_I:%.*]] = shl i64 [[__R_0_I15_I_I]], 3
+// CHECK-NEXT: [[CONV5_I21_I_I:%.*]] = sext i8 [[TMP6]] to i64
// CHECK-NEXT: [[ADD_I22_I_I:%.*]] = add i64 [[MUL_I20_I_I]], -48
// CHECK-NEXT: [[SUB_I23_I_I:%.*]] = add i64 [[ADD_I22_I_I]], [[CONV5_I21_I_I]]
// CHECK-NEXT: [[INCDEC_PTR_I24_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I14_I_I]], i64 1
@@ -1847,9 +1827,30 @@ extern "C" __device__ double test_modf(double x, double* y) {
// CHECK: cleanup.i28.i.i:
// CHECK-NEXT: [[__TAGP_ADDR_1_I26_I_I]] = phi ptr [ [[INCDEC_PTR_I24_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__TAGP_ADDR_0_I14_I_I]], [[WHILE_BODY_I19_I_I]] ]
// CHECK-NEXT: [[__R_1_I27_I_I]] = phi i64 [ [[SUB_I23_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_BODY_I19_I_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[WHILE_COND_I17_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP9]]
+// CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[WHILE_COND_I17_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP6]]
+// CHECK: while.cond.i33.i.i:
+// CHECK-NEXT: [[__TAGP_ADDR_0_I30_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I42_I_I:%.*]], [[CLEANUP_I44_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
+// CHECK-NEXT: [[__R_0_I31_I_I:%.*]] = phi i64 [ [[__R_1_I43_I_I:%.*]], [[CLEANUP_I44_I_I]] ], [ 0, [[ENTRY]] ]
+// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT: [[CMP_NOT_I32_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I32_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I35_I_I:%.*]]
+// CHECK: while.body.i35.i.i:
+// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// CHECK-NEXT: [[OR_COND_I34_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
+// CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[IF_THEN_I41_I_I:%.*]], label [[CLEANUP_I44_I_I]]
+// CHECK: if.then.i41.i.i:
+// CHECK-NEXT: [[MUL_I36_I_I:%.*]] = mul i64 [[__R_0_I31_I_I]], 10
+// CHECK-NEXT: [[CONV5_I37_I_I:%.*]] = sext i8 [[TMP8]] to i64
+// CHECK-NEXT: [[ADD_I38_I_I:%.*]] = add i64 [[MUL_I36_I_I]], -48
+// CHECK-NEXT: [[SUB_I39_I_I:%.*]] = add i64 [[ADD_I38_I_I]], [[CONV5_I37_I_I]]
+// CHECK-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I_I]], i64 1
+// CHECK-NEXT: br label [[CLEANUP_I44_I_I]]
+// CHECK: cleanup.i44.i.i:
+// CHECK-NEXT: [[__TAGP_ADDR_1_I42_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__TAGP_ADDR_0_I30_I_I]], [[WHILE_BODY_I35_I_I]] ]
+// CHECK-NEXT: [[__R_1_I43_I_I]] = phi i64 [ [[SUB_I39_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_BODY_I35_I_I]] ]
+// CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[WHILE_COND_I33_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP9]]
// CHECK: _ZL4nanfPKc.exit:
-// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I44_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_COND_I33_I_I]] ], [ 0, [[CLEANUP_I28_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ]
+// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I28_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ], [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I44_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_COND_I33_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
@@ -1864,26 +1865,26 @@ extern "C" __device__ float test_nanf(const char *tag) {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa [[TBAA3]]
// 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_I17_I_I:%.*]]
+// CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I33_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 [[TBAA3]]
-// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [
-// CHECK-NEXT: i8 120, label [[WHILE_COND_I33_I_I_PREHEADER:%.*]]
-// CHECK-NEXT: i8 88, label [[WHILE_COND_I33_I_I_PREHEADER]]
+// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I17_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: ]
-// CHECK: while.cond.i33.i.i.preheader:
-// CHECK-NEXT: br label [[WHILE_COND_I33_I_I:%.*]]
-// CHECK: while.cond.i33.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I30_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I43_I_I:%.*]], [[CLEANUP_I44_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I33_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[__R_0_I31_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I44_I_I]] ], [ 0, [[WHILE_COND_I33_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: [[CMP_NOT_I32_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I32_I_I]], label [[_ZL3NANPKC_EXIT:%.*]], label [[WHILE_BODY_I35_I_I:%.*]]
-// CHECK: while.body.i35.i.i:
+// 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 [[TBAA3]]
+// 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-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// CHECK-NEXT: [[OR_COND_I34_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
+// 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: if.else.i.i.i:
// CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
// CHECK-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
@@ -1891,54 +1892,33 @@ 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_I44_I_I]]
+// CHECK-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I_I_I]]
// CHECK: if.end31.i.i.i:
-// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I35_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I31_I_I]], 4
+// 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: [[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_I42_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I_I]], i64 1
-// CHECK-NEXT: br label [[CLEANUP_I44_I_I]]
-// CHECK: cleanup.i44.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I43_I_I]] = phi ptr [ [[INCDEC_PTR_I42_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I30_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_I31_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_I33_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP10]]
-// 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 [[TBAA3]]
-// 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_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 [[LOOP6]]
+// 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 [[LOOP10]]
// CHECK: while.cond.i17.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I_I:%.*]], [[CLEANUP_I28_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
-// CHECK-NEXT: [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I27_I_I:%.*]], [[CLEANUP_I28_I_I]] ], [ 0, [[ENTRY]] ]
-// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
+// CHECK-NEXT: [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I_I:%.*]], [[CLEANUP_I28_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
+// CHECK-NEXT: [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I27_I_I:%.*]], [[CLEANUP_I28_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
+// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT: [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
// CHECK-NEXT: br i1 [[CMP_NOT_I16_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I19_I_I:%.*]]
// CHECK: while.body.i19.i.i:
-// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// CHECK-NEXT: [[OR_COND_I18_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
+// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// CHECK-NEXT: [[OR_COND_I18_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
// CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[IF_THEN_I25_I_I:%.*]], label [[CLEANUP_I28_I_I]]
// CHECK: if.then.i25.i.i:
-// CHECK-NEXT: [[MUL_I20_I_I:%.*]] = mul i64 [[__R_0_I15_I_I]], 10
-// CHECK-NEXT: [[CONV5_I21_I_I:%.*]] = sext i8 [[TMP8]] to i64
+// CHECK-NEXT: [[MUL_I20_I_I:%.*]] = shl i64 [[__R_0_I15_I_I]], 3
+// CHECK-NEXT: [[CONV5_I21_I_I:%.*]] = sext i8 [[TMP6]] to i64
// CHECK-NEXT: [[ADD_I22_I_I:%.*]] = add i64 [[MUL_I20_I_I]], -48
// CHECK-NEXT: [[SUB_I23_I_I:%.*]] = add i64 [[ADD_I22_I_I]], [[CONV5_I21_I_I]]
// CHECK-NEXT: [[INCDEC_PTR_I24_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I14_I_I]], i64 1
@@ -1946,9 +1926,30 @@ extern "C" __device__ float test_nanf(const char *tag) {
// CHECK: cleanup.i28.i.i:
// CHECK-NEXT: [[__TAGP_ADDR_1_I26_I_I]] = phi ptr [ [[INCDEC_PTR_I24_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__TAGP_ADDR_0_I14_I_I]], [[WHILE_BODY_I19_I_I]] ]
// CHECK-NEXT: [[__R_1_I27_I_I]] = phi i64 [ [[SUB_I23_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_BODY_I19_I_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[WHILE_COND_I17_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP9]]
+// CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[WHILE_COND_I17_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP6]]
+// CHECK: while.cond.i33.i.i:
+// CHECK-NEXT: [[__TAGP_ADDR_0_I30_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I42_I_I:%.*]], [[CLEANUP_I44_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
+// CHECK-NEXT: [[__R_0_I31_I_I:%.*]] = phi i64 [ [[__R_1_I43_I_I:%.*]], [[CLEANUP_I44_I_I]] ], [ 0, [[ENTRY]] ]
+// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT: [[CMP_NOT_I32_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I32_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I35_I_I:%.*]]
+// CHECK: while.body.i35.i.i:
+// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// CHECK-NEXT: [[OR_COND_I34_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
+// CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[IF_THEN_I41_I_I:%.*]], label [[CLEANUP_I44_I_I]]
+// CHECK: if.then.i41.i.i:
+// CHECK-NEXT: [[MUL_I36_I_I:%.*]] = mul i64 [[__R_0_I31_I_I]], 10
+// CHECK-NEXT: [[CONV5_I37_I_I:%.*]] = sext i8 [[TMP8]] to i64
+// CHECK-NEXT: [[ADD_I38_I_I:%.*]] = add i64 [[MUL_I36_I_I]], -48
+// CHECK-NEXT: [[SUB_I39_I_I:%.*]] = add i64 [[ADD_I38_I_I]], [[CONV5_I37_I_I]]
+// CHECK-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I_I]], i64 1
+// CHECK-NEXT: br label [[CLEANUP_I44_I_I]]
+// CHECK: cleanup.i44.i.i:
+// CHECK-NEXT: [[__TAGP_ADDR_1_I42_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__TAGP_ADDR_0_I30_I_I]], [[WHILE_BODY_I35_I_I]] ]
+// CHECK-NEXT: [[__R_1_I43_I_I]] = phi i64 [ [[SUB_I39_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_BODY_I35_I_I]] ]
+// CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[WHILE_COND_I33_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP9]]
// CHECK: _ZL3nanPKc.exit:
-// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I44_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_COND_I33_I_I]] ], [ 0, [[CLEANUP_I28_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ]
+// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I28_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ], [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I44_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_COND_I33_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
@@ -2160,42 +2161,42 @@ extern "C" __device__ double test_normcdfinv(double x) {
// DEFAULT-LABEL: @test_normf(
// DEFAULT-NEXT: entry:
-// 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-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// DEFAULT: while.body.i:
-// 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 [[TBAA15]]
+// 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 [[TBAA15]]
// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
-// 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: [[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: [[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 [[LOOP19:![0-9]+]]
// DEFAULT: _ZL5normfiPKf.exit:
-// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
+// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_normf(
// FINITEONLY-NEXT: entry:
-// 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-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// FINITEONLY: while.body.i:
-// 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 [[TBAA15]]
+// 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 [[TBAA15]]
// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]]
-// 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: [[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: [[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 [[LOOP19:![0-9]+]]
// FINITEONLY: _ZL5normfiPKf.exit:
-// 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 float @__ocml_sqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
+// 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 float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_normf(int x, const float *y) {
@@ -2204,42 +2205,42 @@ extern "C" __device__ float test_normf(int x, const float *y) {
// DEFAULT-LABEL: @test_norm(
// DEFAULT-NEXT: entry:
-// 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-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// DEFAULT: while.body.i:
-// 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 [[TBAA17]]
+// 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 [[TBAA17]]
// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
-// 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: [[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: [[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 [[LOOP20:![0-9]+]]
// DEFAULT: _ZL4normiPKd.exit:
-// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
+// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_norm(
// FINITEONLY-NEXT: entry:
-// 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-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// FINITEONLY: while.body.i:
-// 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 [[TBAA17]]
+// 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 [[TBAA17]]
// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]]
-// 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: [[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: [[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 [[LOOP20:![0-9]+]]
// FINITEONLY: _ZL4normiPKd.exit:
-// 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 double @__ocml_sqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
+// 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 double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_norm(int x, const double *y) {
@@ -2464,42 +2465,42 @@ extern "C" __device__ double test_rint(double x) {
// DEFAULT-LABEL: @test_rnormf(
// DEFAULT-NEXT: entry:
-// 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-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// DEFAULT: while.body.i:
-// 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 [[TBAA15]]
+// 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 [[TBAA15]]
// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
-// 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: [[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: [[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 [[LOOP21:![0-9]+]]
// DEFAULT: _ZL6rnormfiPKf.exit:
-// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rnormf(
// FINITEONLY-NEXT: entry:
-// 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-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// FINITEONLY: while.body.i:
-// 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 [[TBAA15]]
+// 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 [[TBAA15]]
// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]]
-// 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: [[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: [[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 [[LOOP21:![0-9]+]]
// FINITEONLY: _ZL6rnormfiPKf.exit:
-// 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 float @__ocml_rsqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR14]]
+// 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 float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_rnormf(int x, const float* y) {
@@ -2508,42 +2509,42 @@ extern "C" __device__ float test_rnormf(int x, const float* y) {
// DEFAULT-LABEL: @test_rnorm(
// DEFAULT-NEXT: entry:
-// 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-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// DEFAULT: while.body.i:
-// 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 [[TBAA17]]
+// 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 [[TBAA17]]
// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
-// 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: [[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: [[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 [[LOOP22:![0-9]+]]
// DEFAULT: _ZL5rnormiPKd.exit:
-// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rnorm(
// FINITEONLY-NEXT: entry:
-// 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-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// FINITEONLY: while.body.i:
-// 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 [[TBAA17]]
+// 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 [[TBAA17]]
// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]]
-// 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: [[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: [[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 [[LOOP22:![0-9]+]]
// FINITEONLY: _ZL5rnormiPKd.exit:
-// 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 double @__ocml_rsqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR14]]
+// 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 double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_rnorm(int x, const double* y) {
@@ -3143,30 +3144,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_I20_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL3YNFIF_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
// DEFAULT: for.body.i:
-// 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: [[__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: [[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_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 [[LOOP23:![0-9]+]]
+// 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 [[LOOP23:![0-9]+]]
// DEFAULT: _ZL3ynfif.exit:
-// 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: [[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: ret float [[RETVAL_0_I]]
//
// FINITEONLY-LABEL: @test_ynf(
@@ -3176,30 +3177,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_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
// FINITEONLY: for.body.i:
-// 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: [[__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: [[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_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 [[LOOP23:![0-9]+]]
+// 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 [[LOOP23:![0-9]+]]
// FINITEONLY: _ZL3ynfif.exit:
-// 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: [[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: ret float [[RETVAL_0_I]]
//
extern "C" __device__ float test_ynf(int x, float y) {
@@ -3213,30 +3214,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_I20_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL2YNID_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL2YNID_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
// DEFAULT: for.body.i:
-// 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: [[__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: [[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_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 [[LOOP24:![0-9]+]]
+// 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 [[LOOP24:![0-9]+]]
// DEFAULT: _ZL2ynid.exit:
-// 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: [[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: ret double [[RETVAL_0_I]]
//
// FINITEONLY-LABEL: @test_yn(
@@ -3246,30 +3247,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_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
// FINITEONLY: for.body.i:
-// 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: [[__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: [[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_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 [[LOOP24:![0-9]+]]
+// 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 [[LOOP24:![0-9]+]]
// FINITEONLY: _ZL2ynid.exit:
-// 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: [[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: ret double [[RETVAL_0_I]]
//
extern "C" __device__ double test_yn(int x, double y) {
diff --git a/clang/test/OpenMP/bug57757.cpp b/clang/test/OpenMP/bug57757.cpp
index 7acfe134ddd0b..7894796ac4628 100644
--- a/clang/test/OpenMP/bug57757.cpp
+++ b/clang/test/OpenMP/bug57757.cpp
@@ -32,23 +32,24 @@ 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: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4, !tbaa [[TBAA16:![0-9]+]], !alias.scope !13, !noalias !17
+// 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: 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 [[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: 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: 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 [[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: [[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: 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 9b96c8095b11f..6cc66a0cb1320 100644
--- a/llvm/lib/Passes/PassBuilderPipelines.cpp
+++ b/llvm/lib/Passes/PassBuilderPipelines.cpp
@@ -161,7 +161,7 @@ static cl::opt<bool> EnableModuleInliner("enable-module-inliner",
cl::desc("Enable module inliner"));
static cl::opt<bool> PerformMandatoryInliningsFirst(
- "mandatory-inlining-first", cl::init(false), cl::Hidden,
+ "mandatory-inlining-first", cl::init(true), cl::Hidden,
cl::desc("Perform mandatory inlinings module-wide, before performing "
"inlining"));
@@ -1079,8 +1079,6 @@ 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 5e4d1cdfb3bbf..2a9d3e0df9669 100644
--- a/llvm/test/Other/new-pm-defaults.ll
+++ b/llvm/test/Other/new-pm-defaults.ll
@@ -122,8 +122,6 @@
; 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
@@ -132,12 +130,14 @@
; 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-O-NEXT: Running analysis: AAManager
; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
diff --git a/llvm/test/Other/new-pm-print-pipeline.ll b/llvm/test/Other/new-pm-print-pipeline.ll
index a769bbdc6994e..85029d5cc6d3f 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,{{.*}},instcombine{{.*}}))
+; CHECK-21: require<globals-aa>,function(invalidate<aa>),require<profile-summary>,cgscc(devirt<4>(inline<only-mandatory>,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-defaults.ll b/llvm/test/Other/new-pm-thinlto-defaults.ll
index f040cca5abfd1..b2446fcbd45c1 100644
--- a/llvm/test/Other/new-pm-thinlto-defaults.ll
+++ b/llvm/test/Other/new-pm-thinlto-defaults.ll
@@ -99,8 +99,6 @@
; 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
@@ -109,12 +107,14 @@
; CHECK-O-NEXT: Running pass: InvalidateAnalysisPass<{{.*}}AAManager
; CHECK-O-NEXT: Invalidating analysis: AAManager
; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}ProfileSummaryAnalysis
+; CHECK-PRELINK-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-O-NEXT: Running analysis: AAManager
; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
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 123c1799d44c8..da9fafcc51636 100644
--- a/llvm/test/Other/new-pm-thinlto-postlink-pgo-defaults.ll
+++ b/llvm/test/Other/new-pm-thinlto-postlink-pgo-defaults.ll
@@ -63,7 +63,6 @@
; CHECK-O123-NEXT: Running analysis: BranchProbabilityAnalysis 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
@@ -78,6 +77,7 @@
; 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-O-NEXT: Running analysis: AAManager
; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
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 52e8eeb5460fd..d8a3fcec9026c 100644
--- a/llvm/test/Other/new-pm-thinlto-postlink-samplepgo-defaults.ll
+++ b/llvm/test/Other/new-pm-thinlto-postlink-samplepgo-defaults.ll
@@ -73,7 +73,6 @@
; 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
@@ -87,6 +86,7 @@
; 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-O-NEXT: Running analysis: AAManager
; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
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 22e8162b0fd31..735132de2d86f 100644
--- a/llvm/test/Other/new-pm-thinlto-prelink-pgo-defaults.ll
+++ b/llvm/test/Other/new-pm-thinlto-prelink-pgo-defaults.ll
@@ -87,7 +87,6 @@
; 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
@@ -102,6 +101,7 @@
; 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-O-NEXT: Running analysis: AAManager
; CHECK-O-NEXT: Running analysis: BasicAA
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 ad38a95b115d6..4ae020bb13e3e 100644
--- a/llvm/test/Other/new-pm-thinlto-prelink-samplepgo-defaults.ll
+++ b/llvm/test/Other/new-pm-thinlto-prelink-samplepgo-defaults.ll
@@ -68,7 +68,6 @@
; CHECK-O123-NEXT: Running analysis: BranchProbabilityAnalysis 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
@@ -82,6 +81,7 @@
; 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-O-NEXT: Running analysis: AAManager
; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass
diff --git a/llvm/test/Transforms/Inline/always-inline-phase-ordering.ll b/llvm/test/Transforms/Inline/always-inline-phase-ordering.ll
deleted file mode 100644
index e69ca48344900..0000000000000
--- a/llvm/test/Transforms/Inline/always-inline-phase-ordering.ll
+++ /dev/null
@@ -1,164 +0,0 @@
-; 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}
diff --git a/llvm/test/Transforms/PhaseOrdering/ARM/arm_mult_q15.ll b/llvm/test/Transforms/PhaseOrdering/ARM/arm_mult_q15.ll
index 2aa2b39429d5e..7cbce461a4924 100644
--- a/llvm/test/Transforms/PhaseOrdering/ARM/arm_mult_q15.ll
+++ b/llvm/test/Transforms/PhaseOrdering/ARM/arm_mult_q15.ll
@@ -218,5 +218,5 @@ declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1
attributes #0 = { nounwind "frame-pointer"="all" "min-legal-vector-width"="0" "no-infs-fp-math"="true" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="cortex-m55" "target-features"="+armv8.1-m.main,+dsp,+fp-armv8d16,+fp-armv8d16sp,+fp16,+fp64,+fullfp16,+hwdiv,+lob,+mve,+mve.fp,+ras,+strict-align,+thumb-mode,+vfp2,+vfp2sp,+vfp3d16,+vfp3d16sp,+vfp4d16,+vfp4d16sp,-aes,-bf16,-cdecp0,-cdecp1,-cdecp2,-cdecp3,-cdecp4,-cdecp5,-cdecp6,-cdecp7,-crc,-crypto,-d32,-dotprod,-fp-armv8,-fp-armv8sp,-fp16fml,-hwdiv-arm,-i8mm,-neon,-sb,-sha2,-vfp3,-vfp3sp,-vfp4,-vfp4sp" "unsafe-fp-math"="true" }
attributes #1 = { argmemonly nofree nosync nounwind willreturn }
-attributes #2 = { nounwind "frame-pointer"="all" "min-legal-vector-width"="0" "no-infs-fp-math"="true" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="cortex-m55" "target-features"="+armv8.1-m.main,+dsp,+fp-armv8d16,+fp-armv8d16sp,+fp16,+fp64,+fullfp16,+hwdiv,+lob,+mve,+mve.fp,+ras,+strict-align,+thumb-mode,+vfp2,+vfp2sp,+vfp3d16,+vfp3d16sp,+vfp4d16,+vfp4d16sp,-aes,-bf16,-cdecp0,-cdecp1,-cdecp2,-cdecp3,-cdecp4,-cdecp5,-cdecp6,-cdecp7,-crc,-crypto,-d32,-dotprod,-fp-armv8,-fp-armv8sp,-fp16fml,-hwdiv-arm,-i8mm,-neon,-sb,-sha2,-vfp3,-vfp3sp,-vfp4,-vfp4sp" "unsafe-fp-math"="true" }
+attributes #2 = { alwaysinline nounwind "frame-pointer"="all" "min-legal-vector-width"="0" "no-infs-fp-math"="true" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="cortex-m55" "target-features"="+armv8.1-m.main,+dsp,+fp-armv8d16,+fp-armv8d16sp,+fp16,+fp64,+fullfp16,+hwdiv,+lob,+mve,+mve.fp,+ras,+strict-align,+thumb-mode,+vfp2,+vfp2sp,+vfp3d16,+vfp3d16sp,+vfp4d16,+vfp4d16sp,-aes,-bf16,-cdecp0,-cdecp1,-cdecp2,-cdecp3,-cdecp4,-cdecp5,-cdecp6,-cdecp7,-crc,-crypto,-d32,-dotprod,-fp-armv8,-fp-armv8sp,-fp16fml,-hwdiv-arm,-i8mm,-neon,-sb,-sha2,-vfp3,-vfp3sp,-vfp4,-vfp4sp" "unsafe-fp-math"="true" }
attributes #3 = { nounwind }
More information about the cfe-commits
mailing list