[clang] [llvm] Reapply [SimplifyCFG] Extend jump-threading to allow live local defs (PR #197850)
via cfe-commits
cfe-commits at lists.llvm.org
Fri May 22 12:03:48 PDT 2026
https://github.com/LU-JOHN updated https://github.com/llvm/llvm-project/pull/197850
>From 51ae39ff1cafbe66dd480a788733f278644bcf60 Mon Sep 17 00:00:00 2001
From: John Lu <John.Lu at amd.com>
Date: Thu, 14 May 2026 17:45:28 -0500
Subject: [PATCH 1/2] =?UTF-8?q?Revert=20"Revert=20"[SimplifyCFG]=20Extend?=
=?UTF-8?q?=20jump-threading=20to=20allow=20live=20local=20defs=20?=
=?UTF-8?q?=E2=80=A6=20(#190269)"?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
This reverts commit a76750e6de6aba2223097dc505578556ec245d50.
---
clang/test/CodeGenObjC/exceptions.m | 19 +-
clang/test/Headers/__clang_hip_math.hip | 1383 ++++++++---------
llvm/lib/Transforms/Utils/SimplifyCFG.cpp | 71 +-
.../AArch64/avoid-free-ext-promotion.ll | 36 +-
.../jump-threading-live-on-exit.ll | 195 +++
...hreading-max-jump-threading-live-blocks.ll | 95 ++
6 files changed, 1023 insertions(+), 776 deletions(-)
create mode 100644 llvm/test/Transforms/SimplifyCFG/jump-threading-live-on-exit.ll
create mode 100644 llvm/test/Transforms/SimplifyCFG/jump-threading-max-jump-threading-live-blocks.ll
diff --git a/clang/test/CodeGenObjC/exceptions.m b/clang/test/CodeGenObjC/exceptions.m
index 86d53c1a750cd..66dc0515af03e 100644
--- a/clang/test/CodeGenObjC/exceptions.m
+++ b/clang/test/CodeGenObjC/exceptions.m
@@ -144,18 +144,17 @@ void f4(void) {
// CHECK-NEXT: br label
// -> rethrow
- // finally.call-exit: Predecessors are the @try and @catch fallthroughs
- // as well as the no-match case in the catch mechanism. The i1 is whether
- // to rethrow and should be true only in the last case.
- // CHECK: phi ptr
- // CHECK-NEXT: phi i1
- // CHECK-NEXT: call void @objc_exception_try_exit(ptr nonnull [[EXNDATA]])
+ // finally.call-exit: Predecessor is the no-match case in the catch mechanism
+ // which rethrows.
+ // CHECK: call void @objc_exception_try_exit(ptr nonnull [[EXNDATA]])
// CHECK-NEXT: call void @f4_help(i32 noundef 2)
- // CHECK-NEXT: br i1
- // -> ret, rethrow
+ // CHECK-NEXT: br label
+ // -> rethrow
- // ret:
- // CHECK: ret void
+ // finally.end.critedge: Predecessors are the @try and @catch fallthroughs.
+ // CHECK: call void @objc_exception_try_exit(ptr nonnull [[EXNDATA]])
+ // CHECK-NEXT: call void @f4_help(i32 noundef 2)
+ // CHECK-NEXT: ret void
// Catch mechanism:
// CHECK: call ptr @objc_exception_extract(ptr nonnull [[EXNDATA]])
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index 5a98feb81221e..7c3df21da92a8 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -50,30 +50,27 @@ typedef unsigned long long uint64_t;
// CHECK-LABEL: define dso_local i64 @test___make_mantissa_base8(
// CHECK-SAME: ptr noundef readonly captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] {
// CHECK-NEXT: [[ENTRY:.*]]:
-// CHECK-NEXT: br label %[[WHILE_COND_I:.*]]
-// CHECK: [[WHILE_COND_I]]:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I:%.*]] = phi ptr [ [[P]], %[[ENTRY]] ], [ [[__TAGP_ADDR_1_I:%.*]], %[[CLEANUP_I:.*]] ]
-// CHECK-NEXT: [[__R_0_I:%.*]] = phi i64 [ 0, %[[ENTRY]] ], [ [[__R_1_I:%.*]], %[[CLEANUP_I]] ]
-// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I]], align 1, !tbaa [[CHAR_TBAA7:![0-9]+]]
-// CHECK-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I]], label %[[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT:.*]], label %[[WHILE_BODY_I:.*]]
+// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[P]], align 1, !tbaa [[CHAR_TBAA7:![0-9]+]]
+// CHECK-NEXT: [[CMP_NOT_I1:%.*]] = icmp eq i8 [[TMP0]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I1]], label %[[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT:.*]], label %[[WHILE_BODY_I:.*]]
// CHECK: [[WHILE_BODY_I]]:
-// CHECK-NEXT: [[TMP1:%.*]] = and i8 [[TMP0]], -8
-// CHECK-NEXT: [[OR_COND_I:%.*]] = icmp eq i8 [[TMP1]], 48
-// CHECK-NEXT: br i1 [[OR_COND_I]], label %[[IF_THEN_I:.*]], label %[[CLEANUP_I]]
+// CHECK-NEXT: [[TMP1:%.*]] = phi i8 [ [[TMP3:%.*]], %[[IF_THEN_I:.*]] ], [ [[TMP0]], %[[ENTRY]] ]
+// CHECK-NEXT: [[__R_0_I3:%.*]] = phi i64 [ [[SUB_I:%.*]], %[[IF_THEN_I]] ], [ 0, %[[ENTRY]] ]
+// CHECK-NEXT: [[__TAGP_ADDR_0_I2:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], %[[IF_THEN_I]] ], [ [[P]], %[[ENTRY]] ]
+// CHECK-NEXT: [[TMP2:%.*]] = and i8 [[TMP1]], -8
+// CHECK-NEXT: [[OR_COND_I:%.*]] = icmp eq i8 [[TMP2]], 48
+// CHECK-NEXT: br i1 [[OR_COND_I]], label %[[IF_THEN_I]], label %[[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT]]
// CHECK: [[IF_THEN_I]]:
-// CHECK-NEXT: [[MUL_I:%.*]] = shl i64 [[__R_0_I]], 3
-// CHECK-NEXT: [[CONV5_I:%.*]] = zext nneg i8 [[TMP0]] to i64
+// CHECK-NEXT: [[MUL_I:%.*]] = shl i64 [[__R_0_I3]], 3
+// CHECK-NEXT: [[CONV5_I:%.*]] = zext nneg i8 [[TMP1]] to i64
// CHECK-NEXT: [[ADD_I:%.*]] = add i64 [[MUL_I]], -48
-// CHECK-NEXT: [[SUB_I:%.*]] = add i64 [[ADD_I]], [[CONV5_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I]], i64 1
-// CHECK-NEXT: br label %[[CLEANUP_I]]
-// CHECK: [[CLEANUP_I]]:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I]] = phi ptr [ [[INCDEC_PTR_I]], %[[IF_THEN_I]] ], [ [[__TAGP_ADDR_0_I]], %[[WHILE_BODY_I]] ]
-// CHECK-NEXT: [[__R_1_I]] = phi i64 [ [[SUB_I]], %[[IF_THEN_I]] ], [ [[__R_0_I]], %[[WHILE_BODY_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I]], label %[[WHILE_COND_I]], label %[[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT]], !llvm.loop [[LOOP8:![0-9]+]]
+// CHECK-NEXT: [[SUB_I]] = add i64 [[ADD_I]], [[CONV5_I]]
+// CHECK-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I2]], i64 1
+// CHECK-NEXT: [[TMP3]] = load i8, ptr [[INCDEC_PTR_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// CHECK-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP3]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I]], label %[[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT]], label %[[WHILE_BODY_I]], !llvm.loop [[LOOP8:![0-9]+]]
// CHECK: [[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT]]:
-// CHECK-NEXT: [[RETVAL_2_I:%.*]] = phi i64 [ 0, %[[CLEANUP_I]] ], [ [[__R_0_I]], %[[WHILE_COND_I]] ]
+// CHECK-NEXT: [[RETVAL_2_I:%.*]] = phi i64 [ 0, %[[ENTRY]] ], [ 0, %[[WHILE_BODY_I]] ], [ [[SUB_I]], %[[IF_THEN_I]] ]
// CHECK-NEXT: ret i64 [[RETVAL_2_I]]
//
// AMDGCNSPIRV-LABEL: define spir_func i64 @test___make_mantissa_base8(
@@ -108,30 +105,27 @@ extern "C" __device__ uint64_t test___make_mantissa_base8(const char *p) {
// CHECK-LABEL: define dso_local i64 @test___make_mantissa_base10(
// CHECK-SAME: ptr noundef readonly captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR2]] {
// CHECK-NEXT: [[ENTRY:.*]]:
-// CHECK-NEXT: br label %[[WHILE_COND_I:.*]]
-// CHECK: [[WHILE_COND_I]]:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I:%.*]] = phi ptr [ [[P]], %[[ENTRY]] ], [ [[__TAGP_ADDR_1_I:%.*]], %[[CLEANUP_I:.*]] ]
-// CHECK-NEXT: [[__R_0_I:%.*]] = phi i64 [ 0, %[[ENTRY]] ], [ [[__R_1_I:%.*]], %[[CLEANUP_I]] ]
-// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// CHECK-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I]], label %[[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT:.*]], label %[[WHILE_BODY_I:.*]]
+// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[P]], align 1, !tbaa [[CHAR_TBAA7]]
+// CHECK-NEXT: [[CMP_NOT_I1:%.*]] = icmp eq i8 [[TMP0]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I1]], label %[[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT:.*]], label %[[WHILE_BODY_I:.*]]
// CHECK: [[WHILE_BODY_I]]:
-// CHECK-NEXT: [[TMP1:%.*]] = add i8 [[TMP0]], -48
-// CHECK-NEXT: [[OR_COND_I:%.*]] = icmp ult i8 [[TMP1]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I]], label %[[IF_THEN_I:.*]], label %[[CLEANUP_I]]
+// CHECK-NEXT: [[TMP1:%.*]] = phi i8 [ [[TMP3:%.*]], %[[IF_THEN_I:.*]] ], [ [[TMP0]], %[[ENTRY]] ]
+// CHECK-NEXT: [[__R_0_I3:%.*]] = phi i64 [ [[SUB_I:%.*]], %[[IF_THEN_I]] ], [ 0, %[[ENTRY]] ]
+// CHECK-NEXT: [[__TAGP_ADDR_0_I2:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], %[[IF_THEN_I]] ], [ [[P]], %[[ENTRY]] ]
+// CHECK-NEXT: [[TMP2:%.*]] = add i8 [[TMP1]], -48
+// CHECK-NEXT: [[OR_COND_I:%.*]] = icmp ult i8 [[TMP2]], 10
+// CHECK-NEXT: br i1 [[OR_COND_I]], label %[[IF_THEN_I]], label %[[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT]]
// CHECK: [[IF_THEN_I]]:
-// CHECK-NEXT: [[MUL_I:%.*]] = mul i64 [[__R_0_I]], 10
-// CHECK-NEXT: [[CONV5_I:%.*]] = zext nneg i8 [[TMP0]] to i64
+// CHECK-NEXT: [[MUL_I:%.*]] = mul i64 [[__R_0_I3]], 10
+// CHECK-NEXT: [[CONV5_I:%.*]] = zext nneg i8 [[TMP1]] to i64
// CHECK-NEXT: [[ADD_I:%.*]] = add i64 [[MUL_I]], -48
-// CHECK-NEXT: [[SUB_I:%.*]] = add i64 [[ADD_I]], [[CONV5_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I]], i64 1
-// CHECK-NEXT: br label %[[CLEANUP_I]]
-// CHECK: [[CLEANUP_I]]:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I]] = phi ptr [ [[INCDEC_PTR_I]], %[[IF_THEN_I]] ], [ [[__TAGP_ADDR_0_I]], %[[WHILE_BODY_I]] ]
-// CHECK-NEXT: [[__R_1_I]] = phi i64 [ [[SUB_I]], %[[IF_THEN_I]] ], [ [[__R_0_I]], %[[WHILE_BODY_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I]], label %[[WHILE_COND_I]], label %[[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT]], !llvm.loop [[LOOP11:![0-9]+]]
+// CHECK-NEXT: [[SUB_I]] = add i64 [[ADD_I]], [[CONV5_I]]
+// CHECK-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I2]], i64 1
+// CHECK-NEXT: [[TMP3]] = load i8, ptr [[INCDEC_PTR_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// CHECK-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP3]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I]], label %[[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT]], label %[[WHILE_BODY_I]], !llvm.loop [[LOOP11:![0-9]+]]
// CHECK: [[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT]]:
-// CHECK-NEXT: [[RETVAL_2_I:%.*]] = phi i64 [ 0, %[[CLEANUP_I]] ], [ [[__R_0_I]], %[[WHILE_COND_I]] ]
+// CHECK-NEXT: [[RETVAL_2_I:%.*]] = phi i64 [ 0, %[[ENTRY]] ], [ 0, %[[WHILE_BODY_I]] ], [ [[SUB_I]], %[[IF_THEN_I]] ]
// CHECK-NEXT: ret i64 [[RETVAL_2_I]]
//
// AMDGCNSPIRV-LABEL: define spir_func i64 @test___make_mantissa_base10(
@@ -166,79 +160,71 @@ extern "C" __device__ uint64_t test___make_mantissa_base10(const char *p) {
// CHECK-LABEL: define dso_local i64 @test___make_mantissa_base16(
// CHECK-SAME: ptr noundef readonly captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR2]] {
// CHECK-NEXT: [[ENTRY:.*]]:
-// CHECK-NEXT: br label %[[WHILE_COND_I:.*]]
-// CHECK: [[WHILE_COND_I]]:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I:%.*]] = phi ptr [ [[P]], %[[ENTRY]] ], [ [[__TAGP_ADDR_1_I:%.*]], %[[CLEANUP_I:.*]] ]
-// CHECK-NEXT: [[__R_0_I:%.*]] = phi i64 [ 0, %[[ENTRY]] ], [ [[__R_2_I:%.*]], %[[CLEANUP_I]] ]
-// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// CHECK-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I]], label %[[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT:.*]], label %[[WHILE_BODY_I:.*]]
+// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[P]], align 1, !tbaa [[CHAR_TBAA7]]
+// CHECK-NEXT: [[CMP_NOT_I1:%.*]] = icmp eq i8 [[TMP0]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I1]], label %[[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT:.*]], label %[[WHILE_BODY_I:.*]]
// CHECK: [[WHILE_BODY_I]]:
-// CHECK-NEXT: [[TMP1:%.*]] = add i8 [[TMP0]], -48
-// CHECK-NEXT: [[OR_COND_I:%.*]] = icmp ult i8 [[TMP1]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I]], label %[[IF_END31_I:.*]], label %[[IF_ELSE_I:.*]]
+// CHECK-NEXT: [[TMP1:%.*]] = phi i8 [ [[TMP5:%.*]], %[[IF_END31_I:.*]] ], [ [[TMP0]], %[[ENTRY]] ]
+// CHECK-NEXT: [[__R_0_I3:%.*]] = phi i64 [ [[ADD28_I:%.*]], %[[IF_END31_I]] ], [ 0, %[[ENTRY]] ]
+// CHECK-NEXT: [[__TAGP_ADDR_0_I2:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], %[[IF_END31_I]] ], [ [[P]], %[[ENTRY]] ]
+// CHECK-NEXT: [[TMP2:%.*]] = add i8 [[TMP1]], -48
+// CHECK-NEXT: [[OR_COND_I:%.*]] = icmp ult i8 [[TMP2]], 10
+// CHECK-NEXT: br i1 [[OR_COND_I]], label %[[IF_END31_I]], label %[[IF_ELSE_I:.*]]
// CHECK: [[IF_ELSE_I]]:
-// CHECK-NEXT: [[TMP2:%.*]] = add i8 [[TMP0]], -97
-// CHECK-NEXT: [[OR_COND33_I:%.*]] = icmp ult i8 [[TMP2]], 6
+// CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP1]], -97
+// CHECK-NEXT: [[OR_COND33_I:%.*]] = icmp ult i8 [[TMP3]], 6
// CHECK-NEXT: br i1 [[OR_COND33_I]], label %[[IF_END31_I]], label %[[IF_ELSE17_I:.*]]
// CHECK: [[IF_ELSE17_I]]:
-// CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP0]], -65
-// CHECK-NEXT: [[OR_COND34_I:%.*]] = icmp ult i8 [[TMP3]], 6
-// CHECK-NEXT: br i1 [[OR_COND34_I]], label %[[IF_END31_I]], label %[[CLEANUP_I]]
+// CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP1]], -65
+// CHECK-NEXT: [[OR_COND34_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// CHECK-NEXT: br i1 [[OR_COND34_I]], label %[[IF_END31_I]], label %[[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT]]
// CHECK: [[IF_END31_I]]:
// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I]] ], [ -48, %[[WHILE_BODY_I]] ], [ -55, %[[IF_ELSE17_I]] ]
-// CHECK-NEXT: [[MUL24_I:%.*]] = shl i64 [[__R_0_I]], 4
-// CHECK-NEXT: [[CONV25_I:%.*]] = zext nneg i8 [[TMP0]] to i64
+// CHECK-NEXT: [[MUL24_I:%.*]] = shl i64 [[__R_0_I3]], 4
+// CHECK-NEXT: [[CONV25_I:%.*]] = zext nneg i8 [[TMP1]] to i64
// CHECK-NEXT: [[ADD26_I:%.*]] = add i64 [[MUL24_I]], [[DOTSINK]]
-// CHECK-NEXT: [[ADD28_I:%.*]] = add i64 [[ADD26_I]], [[CONV25_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I]], i64 1
-// CHECK-NEXT: br label %[[CLEANUP_I]]
-// CHECK: [[CLEANUP_I]]:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I]] = phi ptr [ [[INCDEC_PTR_I]], %[[IF_END31_I]] ], [ [[__TAGP_ADDR_0_I]], %[[IF_ELSE17_I]] ]
-// CHECK-NEXT: [[__R_2_I]] = phi i64 [ [[ADD28_I]], %[[IF_END31_I]] ], [ [[__R_0_I]], %[[IF_ELSE17_I]] ]
-// CHECK-NEXT: [[COND_I:%.*]] = phi i1 [ true, %[[IF_END31_I]] ], [ false, %[[IF_ELSE17_I]] ]
-// CHECK-NEXT: br i1 [[COND_I]], label %[[WHILE_COND_I]], label %[[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT]], !llvm.loop [[LOOP12:![0-9]+]]
+// CHECK-NEXT: [[ADD28_I]] = add i64 [[ADD26_I]], [[CONV25_I]]
+// CHECK-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I2]], i64 1
+// CHECK-NEXT: [[TMP5]] = load i8, ptr [[INCDEC_PTR_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// CHECK-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP5]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I]], label %[[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT]], label %[[WHILE_BODY_I]], !llvm.loop [[LOOP12:![0-9]+]]
// CHECK: [[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT]]:
-// CHECK-NEXT: [[RETVAL_2_I:%.*]] = phi i64 [ 0, %[[CLEANUP_I]] ], [ [[__R_0_I]], %[[WHILE_COND_I]] ]
+// CHECK-NEXT: [[RETVAL_2_I:%.*]] = phi i64 [ 0, %[[ENTRY]] ], [ 0, %[[IF_ELSE17_I]] ], [ [[ADD28_I]], %[[IF_END31_I]] ]
// CHECK-NEXT: ret i64 [[RETVAL_2_I]]
//
// AMDGCNSPIRV-LABEL: define spir_func i64 @test___make_mantissa_base16(
// AMDGCNSPIRV-SAME: ptr addrspace(4) noundef readonly captures(none) [[P:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR2]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*]]:
-// AMDGCNSPIRV-NEXT: br label %[[WHILE_COND_I:.*]]
-// AMDGCNSPIRV: [[WHILE_COND_I]]:
-// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_0_I:%.*]] = phi ptr addrspace(4) [ [[P]], %[[ENTRY]] ], [ [[__TAGP_ADDR_1_I:%.*]], %[[CLEANUP_I:.*]] ]
-// AMDGCNSPIRV-NEXT: [[__R_0_I:%.*]] = phi i64 [ 0, %[[ENTRY]] ], [ [[__R_2_I:%.*]], %[[CLEANUP_I]] ]
-// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// AMDGCNSPIRV-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0
-// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I]], label %[[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT:.*]], label %[[WHILE_BODY_I:.*]]
+// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load i8, ptr addrspace(4) [[P]], align 1, !tbaa [[CHAR_TBAA8]]
+// AMDGCNSPIRV-NEXT: [[CMP_NOT_I1:%.*]] = icmp eq i8 [[TMP0]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I1]], label %[[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT:.*]], label %[[WHILE_BODY_I:.*]]
// AMDGCNSPIRV: [[WHILE_BODY_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = add i8 [[TMP0]], -48
-// AMDGCNSPIRV-NEXT: [[OR_COND_I:%.*]] = icmp ult i8 [[TMP1]], 10
-// AMDGCNSPIRV-NEXT: br i1 [[OR_COND_I]], label %[[IF_END31_I:.*]], label %[[IF_ELSE_I:.*]]
+// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = phi i8 [ [[TMP5:%.*]], %[[IF_END31_I:.*]] ], [ [[TMP0]], %[[ENTRY]] ]
+// AMDGCNSPIRV-NEXT: [[__R_0_I3:%.*]] = phi i64 [ [[ADD28_I:%.*]], %[[IF_END31_I]] ], [ 0, %[[ENTRY]] ]
+// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_0_I2:%.*]] = phi ptr addrspace(4) [ [[INCDEC_PTR_I:%.*]], %[[IF_END31_I]] ], [ [[P]], %[[ENTRY]] ]
+// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = add i8 [[TMP1]], -48
+// AMDGCNSPIRV-NEXT: [[OR_COND_I:%.*]] = icmp ult i8 [[TMP2]], 10
+// AMDGCNSPIRV-NEXT: br i1 [[OR_COND_I]], label %[[IF_END31_I]], label %[[IF_ELSE_I:.*]]
// AMDGCNSPIRV: [[IF_ELSE_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = add i8 [[TMP0]], -97
-// AMDGCNSPIRV-NEXT: [[OR_COND33_I:%.*]] = icmp ult i8 [[TMP2]], 6
+// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = add i8 [[TMP1]], -97
+// AMDGCNSPIRV-NEXT: [[OR_COND33_I:%.*]] = icmp ult i8 [[TMP3]], 6
// AMDGCNSPIRV-NEXT: br i1 [[OR_COND33_I]], label %[[IF_END31_I]], label %[[IF_ELSE17_I:.*]]
// AMDGCNSPIRV: [[IF_ELSE17_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = add i8 [[TMP0]], -65
-// AMDGCNSPIRV-NEXT: [[OR_COND34_I:%.*]] = icmp ult i8 [[TMP3]], 6
-// AMDGCNSPIRV-NEXT: br i1 [[OR_COND34_I]], label %[[IF_END31_I]], label %[[CLEANUP_I]]
+// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = add i8 [[TMP1]], -65
+// AMDGCNSPIRV-NEXT: [[OR_COND34_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// AMDGCNSPIRV-NEXT: br i1 [[OR_COND34_I]], label %[[IF_END31_I]], label %[[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT]]
// AMDGCNSPIRV: [[IF_END31_I]]:
// AMDGCNSPIRV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I]] ], [ -48, %[[WHILE_BODY_I]] ], [ -55, %[[IF_ELSE17_I]] ]
-// AMDGCNSPIRV-NEXT: [[MUL24_I:%.*]] = shl i64 [[__R_0_I]], 4
-// AMDGCNSPIRV-NEXT: [[CONV25_I:%.*]] = zext nneg i8 [[TMP0]] to i64
+// AMDGCNSPIRV-NEXT: [[MUL24_I:%.*]] = shl i64 [[__R_0_I3]], 4
+// AMDGCNSPIRV-NEXT: [[CONV25_I:%.*]] = zext nneg i8 [[TMP1]] to i64
// AMDGCNSPIRV-NEXT: [[ADD26_I:%.*]] = add i64 [[MUL24_I]], [[DOTSINK]]
-// AMDGCNSPIRV-NEXT: [[ADD28_I:%.*]] = add i64 [[ADD26_I]], [[CONV25_I]]
-// AMDGCNSPIRV-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__TAGP_ADDR_0_I]], i64 1
-// AMDGCNSPIRV-NEXT: br label %[[CLEANUP_I]]
-// AMDGCNSPIRV: [[CLEANUP_I]]:
-// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_1_I]] = phi ptr addrspace(4) [ [[INCDEC_PTR_I]], %[[IF_END31_I]] ], [ [[__TAGP_ADDR_0_I]], %[[IF_ELSE17_I]] ]
-// AMDGCNSPIRV-NEXT: [[__R_2_I]] = phi i64 [ [[ADD28_I]], %[[IF_END31_I]] ], [ [[__R_0_I]], %[[IF_ELSE17_I]] ]
-// AMDGCNSPIRV-NEXT: [[COND_I:%.*]] = phi i1 [ true, %[[IF_END31_I]] ], [ false, %[[IF_ELSE17_I]] ]
-// AMDGCNSPIRV-NEXT: br i1 [[COND_I]], label %[[WHILE_COND_I]], label %[[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT]], !llvm.loop [[LOOP13:![0-9]+]]
+// AMDGCNSPIRV-NEXT: [[ADD28_I]] = add i64 [[ADD26_I]], [[CONV25_I]]
+// AMDGCNSPIRV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__TAGP_ADDR_0_I2]], i64 1
+// AMDGCNSPIRV-NEXT: [[TMP5]] = load i8, ptr addrspace(4) [[INCDEC_PTR_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// AMDGCNSPIRV-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP5]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I]], label %[[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT]], label %[[WHILE_BODY_I]], !llvm.loop [[LOOP13:![0-9]+]]
// AMDGCNSPIRV: [[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT]]:
-// AMDGCNSPIRV-NEXT: [[RETVAL_2_I:%.*]] = phi i64 [ 0, %[[CLEANUP_I]] ], [ [[__R_0_I]], %[[WHILE_COND_I]] ]
+// AMDGCNSPIRV-NEXT: [[RETVAL_2_I:%.*]] = phi i64 [ 0, %[[ENTRY]] ], [ 0, %[[IF_ELSE17_I]] ], [ [[ADD28_I]], %[[IF_END31_I]] ]
// AMDGCNSPIRV-NEXT: ret i64 [[RETVAL_2_I]]
//
extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
@@ -247,94 +233,85 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
// CHECK-LABEL: define dso_local i64 @test___make_mantissa(
// CHECK-SAME: ptr noundef readonly captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR2]] {
-// CHECK-NEXT: [[ENTRY:.*]]:
+// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[P]], align 1, !tbaa [[CHAR_TBAA7]]
// CHECK-NEXT: [[CMP_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// CHECK-NEXT: br i1 [[CMP_I]], label %[[IF_THEN_I:.*]], label %[[WHILE_COND_I14_I:.*]]
+// CHECK-NEXT: br i1 [[CMP_I]], label %[[IF_THEN_I:.*]], label %[[WHILE_COND_I14_I_PREHEADER:.*]]
+// CHECK: [[WHILE_COND_I14_I_PREHEADER]]:
+// CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[P]], align 1, !tbaa [[CHAR_TBAA7]]
+// CHECK-NEXT: [[CMP_NOT_I17_I5:%.*]] = icmp eq i8 [[TMP1]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I17_I5]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT:.*]], label %[[WHILE_BODY_I18_I:.*]]
// CHECK: [[IF_THEN_I]]:
// CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds nuw i8, ptr [[P]], i64 1
-// CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// CHECK-NEXT: switch i8 [[TMP1]], label %[[WHILE_COND_I_I:.*]] [
-// CHECK-NEXT: i8 120, label %[[WHILE_COND_I30_I_PREHEADER:.*]]
-// CHECK-NEXT: i8 88, label %[[WHILE_COND_I30_I_PREHEADER]]
+// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[INCDEC_PTR_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// CHECK-NEXT: [[CMP_NOT_I_I14:%.*]] = icmp eq i8 [[TMP2]], 0
+// CHECK-NEXT: switch i8 [[TMP2]], label %[[WHILE_COND_I_I_PREHEADER:.*]] [
+// CHECK-NEXT: i8 120, label %[[IF_THEN5_I:.*]]
+// CHECK-NEXT: i8 88, label %[[IF_THEN5_I]]
// CHECK-NEXT: ]
-// CHECK: [[WHILE_COND_I30_I_PREHEADER]]:
-// CHECK-NEXT: br label %[[WHILE_COND_I30_I:.*]]
-// CHECK: [[WHILE_COND_I30_I]]:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I31_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I:%.*]], %[[CLEANUP_I36_I:.*]] ], [ [[INCDEC_PTR_I]], %[[WHILE_COND_I30_I_PREHEADER]] ]
-// CHECK-NEXT: [[__R_0_I32_I:%.*]] = phi i64 [ [[__R_2_I_I:%.*]], %[[CLEANUP_I36_I]] ], [ 0, %[[WHILE_COND_I30_I_PREHEADER]] ]
-// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// CHECK-NEXT: [[CMP_NOT_I33_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I33_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT:.*]], label %[[WHILE_BODY_I34_I:.*]]
-// CHECK: [[WHILE_BODY_I34_I]]:
-// CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// CHECK-NEXT: [[OR_COND_I35_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I35_I]], label %[[IF_END31_I_I:.*]], label %[[IF_ELSE_I_I:.*]]
+// CHECK: [[WHILE_COND_I_I_PREHEADER]]:
+// CHECK-NEXT: br i1 [[CMP_NOT_I_I14]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], label %[[WHILE_BODY_I_I:.*]]
+// CHECK: [[IF_THEN5_I]]:
+// CHECK-NEXT: br i1 [[CMP_NOT_I_I14]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], label %[[WHILE_BODY_I31_I:.*]]
+// CHECK: [[WHILE_BODY_I31_I]]:
+// CHECK-NEXT: [[TMP3:%.*]] = phi i8 [ [[TMP7:%.*]], %[[IF_END31_I_I:.*]] ], [ [[TMP2]], %[[IF_THEN5_I]] ]
+// CHECK-NEXT: [[__R_0_I29_I11:%.*]] = phi i64 [ [[ADD28_I_I:%.*]], %[[IF_END31_I_I]] ], [ 0, %[[IF_THEN5_I]] ]
+// CHECK-NEXT: [[__TAGP_ADDR_0_I28_I10:%.*]] = phi ptr [ [[INCDEC_PTR_I34_I:%.*]], %[[IF_END31_I_I]] ], [ [[INCDEC_PTR_I]], %[[IF_THEN5_I]] ]
+// CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP3]], -48
+// CHECK-NEXT: [[OR_COND_I32_I:%.*]] = icmp ult i8 [[TMP4]], 10
+// CHECK-NEXT: br i1 [[OR_COND_I32_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
+// CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP3]], -97
+// CHECK-NEXT: [[OR_COND33_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
// CHECK-NEXT: br i1 [[OR_COND33_I_I]], label %[[IF_END31_I_I]], label %[[IF_ELSE17_I_I:.*]]
// 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_I36_I]]
+// CHECK-NEXT: [[TMP6:%.*]] = add i8 [[TMP3]], -65
+// CHECK-NEXT: [[OR_COND34_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
+// CHECK-NEXT: br i1 [[OR_COND34_I_I]], label %[[IF_END31_I_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]]
// CHECK: [[IF_END31_I_I]]:
-// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I]] ], [ -48, %[[WHILE_BODY_I34_I]] ], [ -55, %[[IF_ELSE17_I_I]] ]
-// CHECK-NEXT: [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I32_I]], 4
-// CHECK-NEXT: [[CONV25_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I]] ], [ -48, %[[WHILE_BODY_I31_I]] ], [ -55, %[[IF_ELSE17_I_I]] ]
+// CHECK-NEXT: [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I29_I11]], 4
+// CHECK-NEXT: [[CONV25_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// CHECK-NEXT: [[ADD26_I_I:%.*]] = add i64 [[MUL24_I_I]], [[DOTSINK]]
-// CHECK-NEXT: [[ADD28_I_I:%.*]] = add i64 [[ADD26_I_I]], [[CONV25_I_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I40_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I31_I]], i64 1
-// CHECK-NEXT: br label %[[CLEANUP_I36_I]]
-// CHECK: [[CLEANUP_I36_I]]:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I37_I]] = phi ptr [ [[INCDEC_PTR_I40_I]], %[[IF_END31_I_I]] ], [ [[__TAGP_ADDR_0_I31_I]], %[[IF_ELSE17_I_I]] ]
-// CHECK-NEXT: [[__R_2_I_I]] = phi i64 [ [[ADD28_I_I]], %[[IF_END31_I_I]] ], [ [[__R_0_I32_I]], %[[IF_ELSE17_I_I]] ]
-// CHECK-NEXT: [[COND_I_I:%.*]] = phi i1 [ true, %[[IF_END31_I_I]] ], [ false, %[[IF_ELSE17_I_I]] ]
-// CHECK-NEXT: br i1 [[COND_I_I]], label %[[WHILE_COND_I30_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP12]]
-// 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 [[CHAR_TBAA7]]
-// 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-NEXT: [[ADD28_I_I]] = add i64 [[ADD26_I_I]], [[CONV25_I_I]]
+// CHECK-NEXT: [[INCDEC_PTR_I34_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I28_I10]], i64 1
+// CHECK-NEXT: [[TMP7]] = load i8, ptr [[INCDEC_PTR_I34_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// CHECK-NEXT: [[CMP_NOT_I30_I:%.*]] = icmp eq i8 [[TMP7]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I30_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], label %[[WHILE_BODY_I31_I]], !llvm.loop [[LOOP12]]
// 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-NEXT: [[TMP8:%.*]] = phi i8 [ [[TMP10:%.*]], %[[IF_THEN_I_I:.*]] ], [ [[TMP2]], %[[WHILE_COND_I_I_PREHEADER]] ]
+// CHECK-NEXT: [[__R_0_I_I16:%.*]] = phi i64 [ [[SUB_I_I:%.*]], %[[IF_THEN_I_I]] ], [ 0, %[[WHILE_COND_I_I_PREHEADER]] ]
+// CHECK-NEXT: [[__TAGP_ADDR_0_I_I15:%.*]] = phi ptr [ [[INCDEC_PTR_I_I:%.*]], %[[IF_THEN_I_I]] ], [ [[INCDEC_PTR_I]], %[[WHILE_COND_I_I_PREHEADER]] ]
+// CHECK-NEXT: [[TMP9:%.*]] = and i8 [[TMP8]], -8
+// CHECK-NEXT: [[OR_COND_I_I:%.*]] = icmp eq i8 [[TMP9]], 48
+// CHECK-NEXT: br i1 [[OR_COND_I_I]], label %[[IF_THEN_I_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]]
// CHECK: [[IF_THEN_I_I]]:
-// CHECK-NEXT: [[MUL_I_I:%.*]] = shl i64 [[__R_0_I_I]], 3
-// CHECK-NEXT: [[CONV5_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// CHECK-NEXT: [[MUL_I_I:%.*]] = shl i64 [[__R_0_I_I16]], 3
+// CHECK-NEXT: [[CONV5_I_I:%.*]] = zext nneg i8 [[TMP8]] 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 nuw 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 [[LOOP8]]
-// CHECK: [[WHILE_COND_I14_I]]:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I15_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I:%.*]], %[[CLEANUP_I20_I:.*]] ], [ [[P]], %[[ENTRY]] ]
-// CHECK-NEXT: [[__R_0_I16_I:%.*]] = phi i64 [ [[__R_1_I22_I:%.*]], %[[CLEANUP_I20_I]] ], [ 0, %[[ENTRY]] ]
-// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// CHECK-NEXT: [[CMP_NOT_I17_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I17_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], label %[[WHILE_BODY_I18_I:.*]]
+// CHECK-NEXT: [[SUB_I_I]] = add i64 [[ADD_I_I]], [[CONV5_I_I]]
+// CHECK-NEXT: [[INCDEC_PTR_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I_I15]], i64 1
+// CHECK-NEXT: [[TMP10]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// CHECK-NEXT: [[CMP_NOT_I_I:%.*]] = icmp eq i8 [[TMP10]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], label %[[WHILE_BODY_I_I]], !llvm.loop [[LOOP8]]
// CHECK: [[WHILE_BODY_I18_I]]:
-// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// CHECK-NEXT: [[OR_COND_I19_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I19_I]], label %[[IF_THEN_I24_I:.*]], label %[[CLEANUP_I20_I]]
-// CHECK: [[IF_THEN_I24_I]]:
-// CHECK-NEXT: [[MUL_I25_I:%.*]] = mul i64 [[__R_0_I16_I]], 10
-// CHECK-NEXT: [[CONV5_I26_I:%.*]] = zext nneg i8 [[TMP8]] to i64
-// CHECK-NEXT: [[ADD_I27_I:%.*]] = add i64 [[MUL_I25_I]], -48
-// CHECK-NEXT: [[SUB_I28_I:%.*]] = add i64 [[ADD_I27_I]], [[CONV5_I26_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I29_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I]], i64 1
-// CHECK-NEXT: br label %[[CLEANUP_I20_I]]
-// CHECK: [[CLEANUP_I20_I]]:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I21_I]] = phi ptr [ [[INCDEC_PTR_I29_I]], %[[IF_THEN_I24_I]] ], [ [[__TAGP_ADDR_0_I15_I]], %[[WHILE_BODY_I18_I]] ]
-// CHECK-NEXT: [[__R_1_I22_I]] = phi i64 [ [[SUB_I28_I]], %[[IF_THEN_I24_I]] ], [ [[__R_0_I16_I]], %[[WHILE_BODY_I18_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I19_I]], label %[[WHILE_COND_I14_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP11]]
+// CHECK-NEXT: [[TMP11:%.*]] = phi i8 [ [[TMP13:%.*]], %[[IF_THEN_I21_I:.*]] ], [ [[TMP1]], %[[WHILE_COND_I14_I_PREHEADER]] ]
+// CHECK-NEXT: [[__R_0_I16_I7:%.*]] = phi i64 [ [[SUB_I25_I:%.*]], %[[IF_THEN_I21_I]] ], [ 0, %[[WHILE_COND_I14_I_PREHEADER]] ]
+// CHECK-NEXT: [[__TAGP_ADDR_0_I15_I6:%.*]] = phi ptr [ [[INCDEC_PTR_I26_I:%.*]], %[[IF_THEN_I21_I]] ], [ [[P]], %[[WHILE_COND_I14_I_PREHEADER]] ]
+// CHECK-NEXT: [[TMP12:%.*]] = add i8 [[TMP11]], -48
+// CHECK-NEXT: [[OR_COND_I19_I:%.*]] = icmp ult i8 [[TMP12]], 10
+// CHECK-NEXT: br i1 [[OR_COND_I19_I]], label %[[IF_THEN_I21_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]]
+// CHECK: [[IF_THEN_I21_I]]:
+// CHECK-NEXT: [[MUL_I22_I:%.*]] = mul i64 [[__R_0_I16_I7]], 10
+// CHECK-NEXT: [[CONV5_I23_I:%.*]] = zext nneg i8 [[TMP11]] to i64
+// CHECK-NEXT: [[ADD_I24_I:%.*]] = add i64 [[MUL_I22_I]], -48
+// CHECK-NEXT: [[SUB_I25_I]] = add i64 [[ADD_I24_I]], [[CONV5_I23_I]]
+// CHECK-NEXT: [[INCDEC_PTR_I26_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I6]], i64 1
+// CHECK-NEXT: [[TMP13]] = load i8, ptr [[INCDEC_PTR_I26_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// CHECK-NEXT: [[CMP_NOT_I17_I:%.*]] = icmp eq i8 [[TMP13]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I17_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], label %[[WHILE_BODY_I18_I]], !llvm.loop [[LOOP11]]
// CHECK: [[_ZL15__MAKE_MANTISSAPKC_EXIT]]:
-// CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ [[__R_0_I32_I]], %[[WHILE_COND_I30_I]] ], [ [[__R_0_I_I]], %[[WHILE_COND_I_I]] ], [ 0, %[[CLEANUP_I_I]] ], [ 0, %[[CLEANUP_I36_I]] ], [ 0, %[[CLEANUP_I20_I]] ], [ [[__R_0_I16_I]], %[[WHILE_COND_I14_I]] ]
+// CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I]] ], [ 0, %[[WHILE_COND_I14_I_PREHEADER]] ], [ 0, %[[WHILE_COND_I_I_PREHEADER]] ], [ [[SUB_I_I]], %[[IF_THEN_I_I]] ], [ [[ADD28_I_I]], %[[IF_END31_I_I]] ], [ 0, %[[WHILE_BODY_I_I]] ], [ 0, %[[IF_ELSE17_I_I]] ], [ [[SUB_I25_I]], %[[IF_THEN_I21_I]] ], [ 0, %[[WHILE_BODY_I18_I]] ]
// CHECK-NEXT: ret i64 [[RETVAL_0_I]]
//
// AMDGCNSPIRV-LABEL: define spir_func i64 @test___make_mantissa(
@@ -347,53 +324,49 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
// AMDGCNSPIRV-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[P]], i64 1
// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = load i8, ptr addrspace(4) [[INCDEC_PTR_I]], align 1, !tbaa [[CHAR_TBAA8]]
// AMDGCNSPIRV-NEXT: switch i8 [[TMP1]], label %[[WHILE_COND_I_I:.*]] [
-// AMDGCNSPIRV-NEXT: i8 120, label %[[WHILE_COND_I28_I_PREHEADER:.*]]
-// AMDGCNSPIRV-NEXT: i8 88, label %[[WHILE_COND_I28_I_PREHEADER]]
+// AMDGCNSPIRV-NEXT: i8 120, label %[[IF_THEN5_I:.*]]
+// AMDGCNSPIRV-NEXT: i8 88, label %[[IF_THEN5_I]]
// AMDGCNSPIRV-NEXT: ]
-// AMDGCNSPIRV: [[WHILE_COND_I28_I_PREHEADER]]:
-// AMDGCNSPIRV-NEXT: br label %[[WHILE_COND_I28_I:.*]]
-// AMDGCNSPIRV: [[WHILE_COND_I28_I]]:
-// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_0_I29_I:%.*]] = phi ptr addrspace(4) [ [[__TAGP_ADDR_1_I34_I:%.*]], %[[CLEANUP_I_I:.*]] ], [ [[INCDEC_PTR_I]], %[[WHILE_COND_I28_I_PREHEADER]] ]
-// AMDGCNSPIRV-NEXT: [[__R_0_I30_I:%.*]] = phi i64 [ [[__R_2_I_I:%.*]], %[[CLEANUP_I_I]] ], [ 0, %[[WHILE_COND_I28_I_PREHEADER]] ]
-// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I29_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// AMDGCNSPIRV-NEXT: [[CMP_NOT_I31_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I31_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT:.*]], label %[[WHILE_BODY_I32_I:.*]]
+// AMDGCNSPIRV: [[IF_THEN5_I]]:
+// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = load i8, ptr addrspace(4) [[INCDEC_PTR_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// AMDGCNSPIRV-NEXT: [[CMP_NOT_I31_I5:%.*]] = icmp eq i8 [[TMP2]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I31_I5]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT:.*]], label %[[WHILE_BODY_I32_I:.*]]
// AMDGCNSPIRV: [[WHILE_BODY_I32_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// AMDGCNSPIRV-NEXT: [[OR_COND_I33_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// AMDGCNSPIRV-NEXT: br i1 [[OR_COND_I33_I]], label %[[IF_END31_I_I:.*]], label %[[IF_ELSE_I_I:.*]]
+// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = phi i8 [ [[TMP7:%.*]], %[[IF_END31_I_I:.*]] ], [ [[TMP2]], %[[IF_THEN5_I]] ]
+// AMDGCNSPIRV-NEXT: [[__R_0_I30_I7:%.*]] = phi i64 [ [[ADD28_I_I:%.*]], %[[IF_END31_I_I]] ], [ 0, %[[IF_THEN5_I]] ]
+// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_0_I29_I6:%.*]] = phi ptr addrspace(4) [ [[INCDEC_PTR_I36_I:%.*]], %[[IF_END31_I_I]] ], [ [[INCDEC_PTR_I]], %[[IF_THEN5_I]] ]
+// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = add i8 [[TMP3]], -48
+// AMDGCNSPIRV-NEXT: [[OR_COND_I33_I:%.*]] = icmp ult i8 [[TMP4]], 10
+// AMDGCNSPIRV-NEXT: br i1 [[OR_COND_I33_I]], label %[[IF_END31_I_I]], label %[[IF_ELSE_I_I:.*]]
// AMDGCNSPIRV: [[IF_ELSE_I_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
-// AMDGCNSPIRV-NEXT: [[OR_COND33_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = add i8 [[TMP3]], -97
+// AMDGCNSPIRV-NEXT: [[OR_COND33_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
// AMDGCNSPIRV-NEXT: br i1 [[OR_COND33_I_I]], label %[[IF_END31_I_I]], label %[[IF_ELSE17_I_I:.*]]
// AMDGCNSPIRV: [[IF_ELSE17_I_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
-// AMDGCNSPIRV-NEXT: [[OR_COND34_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// AMDGCNSPIRV-NEXT: br i1 [[OR_COND34_I_I]], label %[[IF_END31_I_I]], label %[[CLEANUP_I_I]]
+// AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = add i8 [[TMP3]], -65
+// AMDGCNSPIRV-NEXT: [[OR_COND34_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
+// AMDGCNSPIRV-NEXT: br i1 [[OR_COND34_I_I]], label %[[IF_END31_I_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]]
// AMDGCNSPIRV: [[IF_END31_I_I]]:
// AMDGCNSPIRV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I]] ], [ -48, %[[WHILE_BODY_I32_I]] ], [ -55, %[[IF_ELSE17_I_I]] ]
-// AMDGCNSPIRV-NEXT: [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I30_I]], 4
-// AMDGCNSPIRV-NEXT: [[CONV25_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// AMDGCNSPIRV-NEXT: [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I30_I7]], 4
+// AMDGCNSPIRV-NEXT: [[CONV25_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// AMDGCNSPIRV-NEXT: [[ADD26_I_I:%.*]] = add i64 [[MUL24_I_I]], [[DOTSINK]]
-// AMDGCNSPIRV-NEXT: [[ADD28_I_I:%.*]] = add i64 [[ADD26_I_I]], [[CONV25_I_I]]
-// AMDGCNSPIRV-NEXT: [[INCDEC_PTR_I37_I:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__TAGP_ADDR_0_I29_I]], i64 1
-// AMDGCNSPIRV-NEXT: br label %[[CLEANUP_I_I]]
-// AMDGCNSPIRV: [[CLEANUP_I_I]]:
-// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_1_I34_I]] = phi ptr addrspace(4) [ [[INCDEC_PTR_I37_I]], %[[IF_END31_I_I]] ], [ [[__TAGP_ADDR_0_I29_I]], %[[IF_ELSE17_I_I]] ]
-// AMDGCNSPIRV-NEXT: [[__R_2_I_I]] = phi i64 [ [[ADD28_I_I]], %[[IF_END31_I_I]] ], [ [[__R_0_I30_I]], %[[IF_ELSE17_I_I]] ]
-// AMDGCNSPIRV-NEXT: [[COND_I_I:%.*]] = phi i1 [ true, %[[IF_END31_I_I]] ], [ false, %[[IF_ELSE17_I_I]] ]
-// AMDGCNSPIRV-NEXT: br i1 [[COND_I_I]], label %[[WHILE_COND_I28_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP13]]
+// AMDGCNSPIRV-NEXT: [[ADD28_I_I]] = add i64 [[ADD26_I_I]], [[CONV25_I_I]]
+// AMDGCNSPIRV-NEXT: [[INCDEC_PTR_I36_I]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__TAGP_ADDR_0_I29_I6]], i64 1
+// AMDGCNSPIRV-NEXT: [[TMP7]] = load i8, ptr addrspace(4) [[INCDEC_PTR_I36_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// AMDGCNSPIRV-NEXT: [[CMP_NOT_I31_I:%.*]] = icmp eq i8 [[TMP7]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I31_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], label %[[WHILE_BODY_I32_I]], !llvm.loop [[LOOP13]]
// AMDGCNSPIRV: [[WHILE_COND_I_I]]:
// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_0_I_I:%.*]] = phi ptr addrspace(4) [ [[__TAGP_ADDR_1_I_I:%.*]], %[[WHILE_BODY_I_I:.*]] ], [ [[INCDEC_PTR_I]], %[[IF_THEN_I]] ]
// AMDGCNSPIRV-NEXT: [[__R_0_I_I:%.*]] = phi i64 [ [[__R_1_I_I:%.*]], %[[WHILE_BODY_I_I]] ], [ 0, %[[IF_THEN_I]] ]
-// AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// AMDGCNSPIRV-NEXT: [[CMP_NOT_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// AMDGCNSPIRV-NEXT: [[CMP_NOT_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], label %[[WHILE_BODY_I_I]]
// AMDGCNSPIRV: [[WHILE_BODY_I_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// AMDGCNSPIRV-NEXT: [[OR_COND_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = and i8 [[TMP8]], -8
+// AMDGCNSPIRV-NEXT: [[OR_COND_I_I:%.*]] = icmp eq i8 [[TMP9]], 48
// AMDGCNSPIRV-NEXT: [[MUL_I_I:%.*]] = shl i64 [[__R_0_I_I]], 3
-// AMDGCNSPIRV-NEXT: [[CONV5_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// AMDGCNSPIRV-NEXT: [[CONV5_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
// AMDGCNSPIRV-NEXT: [[ADD_I_I:%.*]] = add i64 [[MUL_I_I]], -48
// AMDGCNSPIRV-NEXT: [[SUB_I_I:%.*]] = add i64 [[ADD_I_I]], [[CONV5_I_I]]
// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_1_I_I_IDX:%.*]] = zext i1 [[OR_COND_I_I]] to i64
@@ -403,14 +376,14 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
// AMDGCNSPIRV: [[WHILE_COND_I14_I]]:
// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_0_I15_I:%.*]] = phi ptr addrspace(4) [ [[__TAGP_ADDR_1_I25_I:%.*]], %[[WHILE_BODY_I18_I:.*]] ], [ [[P]], %[[ENTRY]] ]
// AMDGCNSPIRV-NEXT: [[__R_0_I16_I:%.*]] = phi i64 [ [[__R_1_I26_I:%.*]], %[[WHILE_BODY_I18_I]] ], [ 0, %[[ENTRY]] ]
-// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I15_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// AMDGCNSPIRV-NEXT: [[CMP_NOT_I17_I:%.*]] = icmp eq i8 [[TMP8]], 0
+// AMDGCNSPIRV-NEXT: [[TMP10:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I15_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// AMDGCNSPIRV-NEXT: [[CMP_NOT_I17_I:%.*]] = icmp eq i8 [[TMP10]], 0
// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I17_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], label %[[WHILE_BODY_I18_I]]
// AMDGCNSPIRV: [[WHILE_BODY_I18_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// AMDGCNSPIRV-NEXT: [[OR_COND_I19_I:%.*]] = icmp ult i8 [[TMP9]], 10
+// AMDGCNSPIRV-NEXT: [[TMP11:%.*]] = add i8 [[TMP10]], -48
+// AMDGCNSPIRV-NEXT: [[OR_COND_I19_I:%.*]] = icmp ult i8 [[TMP11]], 10
// AMDGCNSPIRV-NEXT: [[MUL_I20_I:%.*]] = mul i64 [[__R_0_I16_I]], 10
-// AMDGCNSPIRV-NEXT: [[CONV5_I21_I:%.*]] = zext nneg i8 [[TMP8]] to i64
+// AMDGCNSPIRV-NEXT: [[CONV5_I21_I:%.*]] = zext nneg i8 [[TMP10]] to i64
// AMDGCNSPIRV-NEXT: [[ADD_I22_I:%.*]] = add i64 [[MUL_I20_I]], -48
// AMDGCNSPIRV-NEXT: [[SUB_I23_I:%.*]] = add i64 [[ADD_I22_I]], [[CONV5_I21_I]]
// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_1_I25_I_IDX:%.*]] = zext i1 [[OR_COND_I19_I]] to i64
@@ -418,7 +391,7 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
// AMDGCNSPIRV-NEXT: [[__R_1_I26_I]] = select i1 [[OR_COND_I19_I]], i64 [[SUB_I23_I]], i64 [[__R_0_I16_I]]
// AMDGCNSPIRV-NEXT: br i1 [[OR_COND_I19_I]], label %[[WHILE_COND_I14_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP12]]
// AMDGCNSPIRV: [[_ZL15__MAKE_MANTISSAPKC_EXIT]]:
-// AMDGCNSPIRV-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ [[__R_0_I30_I]], %[[WHILE_COND_I28_I]] ], [ [[__R_0_I_I]], %[[WHILE_COND_I_I]] ], [ 0, %[[WHILE_BODY_I_I]] ], [ 0, %[[CLEANUP_I_I]] ], [ 0, %[[WHILE_BODY_I18_I]] ], [ [[__R_0_I16_I]], %[[WHILE_COND_I14_I]] ]
+// AMDGCNSPIRV-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I]] ], [ [[ADD28_I_I]], %[[IF_END31_I_I]] ], [ [[__R_0_I_I]], %[[WHILE_COND_I_I]] ], [ 0, %[[WHILE_BODY_I_I]] ], [ 0, %[[IF_ELSE17_I_I]] ], [ 0, %[[WHILE_BODY_I18_I]] ], [ [[__R_0_I16_I]], %[[WHILE_COND_I14_I]] ]
// AMDGCNSPIRV-NEXT: ret i64 [[RETVAL_0_I]]
//
extern "C" __device__ uint64_t test___make_mantissa(const char *p) {
@@ -4349,99 +4322,90 @@ extern "C" __device__ double test_modf(double x, double* y) {
// DEFAULT-LABEL: define dso_local float @test_nanf(
// DEFAULT-SAME: ptr noundef readonly captures(none) [[TAG:%.*]]) local_unnamed_addr #[[ATTR2]] {
-// DEFAULT-NEXT: [[ENTRY:.*]]:
+// DEFAULT-NEXT: [[ENTRY:.*:]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG]], align 1, !tbaa [[CHAR_TBAA7]]
// DEFAULT-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// DEFAULT-NEXT: br i1 [[CMP_I_I]], label %[[IF_THEN_I_I:.*]], label %[[WHILE_COND_I14_I_I:.*]]
+// DEFAULT-NEXT: br i1 [[CMP_I_I]], label %[[IF_THEN_I_I:.*]], label %[[WHILE_COND_I14_I_I_PREHEADER:.*]]
+// DEFAULT: [[WHILE_COND_I14_I_I_PREHEADER]]:
+// DEFAULT-NEXT: [[TMP1:%.*]] = load i8, ptr [[TAG]], align 1, !tbaa [[CHAR_TBAA7]]
+// DEFAULT-NEXT: [[CMP_NOT_I17_I_I5:%.*]] = icmp eq i8 [[TMP1]], 0
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I17_I_I5]], label %[[_ZL4NANFPKC_EXIT:.*]], label %[[WHILE_BODY_I18_I_I:.*]]
// DEFAULT: [[IF_THEN_I_I]]:
// DEFAULT-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[TAG]], i64 1
-// DEFAULT-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// DEFAULT-NEXT: switch i8 [[TMP1]], label %[[WHILE_COND_I_I_I:.*]] [
-// DEFAULT-NEXT: i8 120, label %[[WHILE_COND_I30_I_I_PREHEADER:.*]]
-// DEFAULT-NEXT: i8 88, label %[[WHILE_COND_I30_I_I_PREHEADER]]
+// DEFAULT-NEXT: [[TMP2:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// DEFAULT-NEXT: [[CMP_NOT_I_I_I14:%.*]] = icmp eq i8 [[TMP2]], 0
+// DEFAULT-NEXT: switch i8 [[TMP2]], label %[[WHILE_COND_I_I_I_PREHEADER:.*]] [
+// DEFAULT-NEXT: i8 120, label %[[IF_THEN5_I_I:.*]]
+// DEFAULT-NEXT: i8 88, label %[[IF_THEN5_I_I]]
// DEFAULT-NEXT: ]
-// DEFAULT: [[WHILE_COND_I30_I_I_PREHEADER]]:
-// DEFAULT-NEXT: br label %[[WHILE_COND_I30_I_I:.*]]
-// DEFAULT: [[WHILE_COND_I30_I_I]]:
-// DEFAULT-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I_I:%.*]], %[[CLEANUP_I36_I_I:.*]] ], [ [[INCDEC_PTR_I_I]], %[[WHILE_COND_I30_I_I_PREHEADER]] ]
-// DEFAULT-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], %[[CLEANUP_I36_I_I]] ], [ 0, %[[WHILE_COND_I30_I_I_PREHEADER]] ]
-// DEFAULT-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// DEFAULT-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// DEFAULT-NEXT: br i1 [[CMP_NOT_I33_I_I]], label %[[_ZL4NANFPKC_EXIT:.*]], label %[[WHILE_BODY_I34_I_I:.*]]
-// DEFAULT: [[WHILE_BODY_I34_I_I]]:
-// DEFAULT-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// DEFAULT-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// DEFAULT-NEXT: br i1 [[OR_COND_I35_I_I]], label %[[IF_END31_I_I_I:.*]], label %[[IF_ELSE_I_I_I:.*]]
+// DEFAULT: [[WHILE_COND_I_I_I_PREHEADER]]:
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I_I_I14]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I_I_I:.*]]
+// DEFAULT: [[IF_THEN5_I_I]]:
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I_I_I14]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I31_I_I:.*]]
+// DEFAULT: [[WHILE_BODY_I31_I_I]]:
+// DEFAULT-NEXT: [[TMP3:%.*]] = phi i8 [ [[TMP7:%.*]], %[[IF_END31_I_I_I:.*]] ], [ [[TMP2]], %[[IF_THEN5_I_I]] ]
+// DEFAULT-NEXT: [[__R_0_I29_I_I11:%.*]] = phi i64 [ [[ADD28_I_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ 0, %[[IF_THEN5_I_I]] ]
+// DEFAULT-NEXT: [[__TAGP_ADDR_0_I28_I_I10:%.*]] = phi ptr [ [[INCDEC_PTR_I34_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ [[INCDEC_PTR_I_I]], %[[IF_THEN5_I_I]] ]
+// DEFAULT-NEXT: [[TMP4:%.*]] = add i8 [[TMP3]], -48
+// DEFAULT-NEXT: [[OR_COND_I32_I_I:%.*]] = icmp ult i8 [[TMP4]], 10
+// DEFAULT-NEXT: br i1 [[OR_COND_I32_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE_I_I_I:.*]]
// DEFAULT: [[IF_ELSE_I_I_I]]:
-// DEFAULT-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
-// DEFAULT-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// DEFAULT-NEXT: [[TMP5:%.*]] = add i8 [[TMP3]], -97
+// DEFAULT-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
// DEFAULT-NEXT: br i1 [[OR_COND33_I_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE17_I_I_I:.*]]
// DEFAULT: [[IF_ELSE17_I_I_I]]:
-// DEFAULT-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
-// DEFAULT-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// DEFAULT-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[CLEANUP_I36_I_I]]
+// DEFAULT-NEXT: [[TMP6:%.*]] = add i8 [[TMP3]], -65
+// DEFAULT-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
+// DEFAULT-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[_ZL4NANFPKC_EXIT]]
// DEFAULT: [[IF_END31_I_I_I]]:
-// DEFAULT-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I34_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
-// DEFAULT-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4
-// DEFAULT-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// DEFAULT-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
+// DEFAULT-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I29_I_I11]], 4
+// DEFAULT-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// DEFAULT-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
-// DEFAULT-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
-// DEFAULT-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I31_I_I]], i64 1
-// DEFAULT-NEXT: br label %[[CLEANUP_I36_I_I]]
-// DEFAULT: [[CLEANUP_I36_I_I]]:
-// DEFAULT-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], %[[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// DEFAULT-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// DEFAULT-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, %[[IF_END31_I_I_I]] ], [ false, %[[IF_ELSE17_I_I_I]] ]
-// DEFAULT-NEXT: br i1 [[COND_I_I_I]], label %[[WHILE_COND_I30_I_I]], label %[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP12]]
-// DEFAULT: [[WHILE_COND_I_I_I]]:
-// DEFAULT-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], %[[CLEANUP_I_I_I:.*]] ], [ [[INCDEC_PTR_I_I]], %[[IF_THEN_I_I]] ]
-// DEFAULT-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], %[[CLEANUP_I_I_I]] ], [ 0, %[[IF_THEN_I_I]] ]
-// DEFAULT-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// DEFAULT-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
-// DEFAULT-NEXT: br i1 [[CMP_NOT_I_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I_I_I:.*]]
+// DEFAULT-NEXT: [[ADD28_I_I_I]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I34_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I28_I_I10]], i64 1
+// DEFAULT-NEXT: [[TMP7]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// DEFAULT-NEXT: [[CMP_NOT_I30_I_I:%.*]] = icmp eq i8 [[TMP7]], 0
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I30_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I31_I_I]], !llvm.loop [[LOOP12]]
// DEFAULT: [[WHILE_BODY_I_I_I]]:
-// DEFAULT-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// DEFAULT-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
-// DEFAULT-NEXT: br i1 [[OR_COND_I_I_I]], label %[[IF_THEN_I_I_I:.*]], label %[[CLEANUP_I_I_I]]
+// DEFAULT-NEXT: [[TMP8:%.*]] = phi i8 [ [[TMP10:%.*]], %[[IF_THEN_I_I_I:.*]] ], [ [[TMP2]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[__R_0_I_I_I16:%.*]] = phi i64 [ [[SUB_I_I_I:%.*]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[__TAGP_ADDR_0_I_I_I15:%.*]] = phi ptr [ [[INCDEC_PTR_I_I_I:%.*]], %[[IF_THEN_I_I_I]] ], [ [[INCDEC_PTR_I_I]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[TMP9:%.*]] = and i8 [[TMP8]], -8
+// DEFAULT-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP9]], 48
+// DEFAULT-NEXT: br i1 [[OR_COND_I_I_I]], label %[[IF_THEN_I_I_I]], label %[[_ZL4NANFPKC_EXIT]]
// DEFAULT: [[IF_THEN_I_I_I]]:
-// DEFAULT-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
-// DEFAULT-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// DEFAULT-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I16]], 3
+// DEFAULT-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
// DEFAULT-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
-// DEFAULT-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
-// DEFAULT-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1
-// DEFAULT-NEXT: br label %[[CLEANUP_I_I_I]]
-// DEFAULT: [[CLEANUP_I_I_I]]:
-// DEFAULT-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], %[[WHILE_BODY_I_I_I]] ]
-// DEFAULT-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_BODY_I_I_I]] ]
-// DEFAULT-NEXT: br i1 [[OR_COND_I_I_I]], label %[[WHILE_COND_I_I_I]], label %[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP8]]
-// DEFAULT: [[WHILE_COND_I14_I_I]]:
-// DEFAULT-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I_I:%.*]], %[[CLEANUP_I20_I_I:.*]] ], [ [[TAG]], %[[ENTRY]] ]
-// DEFAULT-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]], %[[CLEANUP_I20_I_I]] ], [ 0, %[[ENTRY]] ]
-// DEFAULT-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// DEFAULT-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// DEFAULT-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I18_I_I:.*]]
+// DEFAULT-NEXT: [[SUB_I_I_I]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I_I_I15]], i64 1
+// DEFAULT-NEXT: [[TMP10]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// DEFAULT-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP10]], 0
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I_I_I]], !llvm.loop [[LOOP8]]
// DEFAULT: [[WHILE_BODY_I18_I_I]]:
-// DEFAULT-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// DEFAULT-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// DEFAULT-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[IF_THEN_I24_I_I:.*]], label %[[CLEANUP_I20_I_I]]
-// DEFAULT: [[IF_THEN_I24_I_I]]:
-// DEFAULT-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
-// DEFAULT-NEXT: [[CONV5_I26_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
-// DEFAULT-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48
-// DEFAULT-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]], [[CONV5_I26_I_I]]
-// DEFAULT-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I_I]], i64 1
-// DEFAULT-NEXT: br label %[[CLEANUP_I20_I_I]]
-// DEFAULT: [[CLEANUP_I20_I_I]]:
-// DEFAULT-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [ [[INCDEC_PTR_I29_I_I]], %[[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]], %[[WHILE_BODY_I18_I_I]] ]
-// DEFAULT-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]], %[[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_BODY_I18_I_I]] ]
-// DEFAULT-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[WHILE_COND_I14_I_I]], label %[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP11]]
+// DEFAULT-NEXT: [[TMP11:%.*]] = phi i8 [ [[TMP13:%.*]], %[[IF_THEN_I21_I_I:.*]] ], [ [[TMP1]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[__R_0_I16_I_I7:%.*]] = phi i64 [ [[SUB_I25_I_I:%.*]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[__TAGP_ADDR_0_I15_I_I6:%.*]] = phi ptr [ [[INCDEC_PTR_I26_I_I:%.*]], %[[IF_THEN_I21_I_I]] ], [ [[TAG]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[TMP12:%.*]] = add i8 [[TMP11]], -48
+// DEFAULT-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP12]], 10
+// DEFAULT-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[IF_THEN_I21_I_I]], label %[[_ZL4NANFPKC_EXIT]]
+// DEFAULT: [[IF_THEN_I21_I_I]]:
+// DEFAULT-NEXT: [[MUL_I22_I_I:%.*]] = mul i64 [[__R_0_I16_I_I7]], 10
+// DEFAULT-NEXT: [[CONV5_I23_I_I:%.*]] = zext nneg i8 [[TMP11]] to i64
+// DEFAULT-NEXT: [[ADD_I24_I_I:%.*]] = add i64 [[MUL_I22_I_I]], -48
+// DEFAULT-NEXT: [[SUB_I25_I_I]] = add i64 [[ADD_I24_I_I]], [[CONV5_I23_I_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I26_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I_I6]], i64 1
+// DEFAULT-NEXT: [[TMP13]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// DEFAULT-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP13]], 0
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I18_I_I]], !llvm.loop [[LOOP11]]
// DEFAULT: [[_ZL4NANFPKC_EXIT]]:
-// DEFAULT-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ [[__R_0_I32_I_I]], %[[WHILE_COND_I30_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_COND_I_I_I]] ], [ 0, %[[CLEANUP_I_I_I]] ], [ 0, %[[CLEANUP_I36_I_I]] ], [ 0, %[[CLEANUP_I20_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_COND_I14_I_I]] ]
+// DEFAULT-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
// DEFAULT-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32
// DEFAULT-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303
// DEFAULT-NEXT: [[BF_SET9_I:%.*]] = or disjoint i32 [[BF_VALUE_I]], 2143289344
-// DEFAULT-NEXT: [[TMP10:%.*]] = bitcast i32 [[BF_SET9_I]] to float
-// DEFAULT-NEXT: ret float [[TMP10]]
+// DEFAULT-NEXT: [[TMP14:%.*]] = bitcast i32 [[BF_SET9_I]] to float
+// DEFAULT-NEXT: ret float [[TMP14]]
//
// FINITEONLY-LABEL: define dso_local nofpclass(nan inf) float @test_nanf(
// FINITEONLY-SAME: ptr noundef readonly captures(none) [[TAG:%.*]]) local_unnamed_addr #[[ATTR3]] {
@@ -4450,195 +4414,177 @@ extern "C" __device__ double test_modf(double x, double* y) {
//
// APPROX-LABEL: define dso_local float @test_nanf(
// APPROX-SAME: ptr noundef readonly captures(none) [[TAG:%.*]]) local_unnamed_addr #[[ATTR2]] {
-// APPROX-NEXT: [[ENTRY:.*]]:
+// APPROX-NEXT: [[ENTRY:.*:]]
// APPROX-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG]], align 1, !tbaa [[CHAR_TBAA7]]
// APPROX-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// APPROX-NEXT: br i1 [[CMP_I_I]], label %[[IF_THEN_I_I:.*]], label %[[WHILE_COND_I14_I_I:.*]]
+// APPROX-NEXT: br i1 [[CMP_I_I]], label %[[IF_THEN_I_I:.*]], label %[[WHILE_COND_I14_I_I_PREHEADER:.*]]
+// APPROX: [[WHILE_COND_I14_I_I_PREHEADER]]:
+// APPROX-NEXT: [[TMP1:%.*]] = load i8, ptr [[TAG]], align 1, !tbaa [[CHAR_TBAA7]]
+// APPROX-NEXT: [[CMP_NOT_I17_I_I5:%.*]] = icmp eq i8 [[TMP1]], 0
+// APPROX-NEXT: br i1 [[CMP_NOT_I17_I_I5]], label %[[_ZL4NANFPKC_EXIT:.*]], label %[[WHILE_BODY_I18_I_I:.*]]
// APPROX: [[IF_THEN_I_I]]:
// APPROX-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[TAG]], i64 1
-// APPROX-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// APPROX-NEXT: switch i8 [[TMP1]], label %[[WHILE_COND_I_I_I:.*]] [
-// APPROX-NEXT: i8 120, label %[[WHILE_COND_I30_I_I_PREHEADER:.*]]
-// APPROX-NEXT: i8 88, label %[[WHILE_COND_I30_I_I_PREHEADER]]
+// APPROX-NEXT: [[TMP2:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// APPROX-NEXT: [[CMP_NOT_I_I_I14:%.*]] = icmp eq i8 [[TMP2]], 0
+// APPROX-NEXT: switch i8 [[TMP2]], label %[[WHILE_COND_I_I_I_PREHEADER:.*]] [
+// APPROX-NEXT: i8 120, label %[[IF_THEN5_I_I:.*]]
+// APPROX-NEXT: i8 88, label %[[IF_THEN5_I_I]]
// APPROX-NEXT: ]
-// APPROX: [[WHILE_COND_I30_I_I_PREHEADER]]:
-// APPROX-NEXT: br label %[[WHILE_COND_I30_I_I:.*]]
-// APPROX: [[WHILE_COND_I30_I_I]]:
-// APPROX-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I_I:%.*]], %[[CLEANUP_I36_I_I:.*]] ], [ [[INCDEC_PTR_I_I]], %[[WHILE_COND_I30_I_I_PREHEADER]] ]
-// APPROX-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], %[[CLEANUP_I36_I_I]] ], [ 0, %[[WHILE_COND_I30_I_I_PREHEADER]] ]
-// APPROX-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// APPROX-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// APPROX-NEXT: br i1 [[CMP_NOT_I33_I_I]], label %[[_ZL4NANFPKC_EXIT:.*]], label %[[WHILE_BODY_I34_I_I:.*]]
-// APPROX: [[WHILE_BODY_I34_I_I]]:
-// APPROX-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// APPROX-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// APPROX-NEXT: br i1 [[OR_COND_I35_I_I]], label %[[IF_END31_I_I_I:.*]], label %[[IF_ELSE_I_I_I:.*]]
+// APPROX: [[WHILE_COND_I_I_I_PREHEADER]]:
+// APPROX-NEXT: br i1 [[CMP_NOT_I_I_I14]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I_I_I:.*]]
+// APPROX: [[IF_THEN5_I_I]]:
+// APPROX-NEXT: br i1 [[CMP_NOT_I_I_I14]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I31_I_I:.*]]
+// APPROX: [[WHILE_BODY_I31_I_I]]:
+// APPROX-NEXT: [[TMP3:%.*]] = phi i8 [ [[TMP7:%.*]], %[[IF_END31_I_I_I:.*]] ], [ [[TMP2]], %[[IF_THEN5_I_I]] ]
+// APPROX-NEXT: [[__R_0_I29_I_I11:%.*]] = phi i64 [ [[ADD28_I_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ 0, %[[IF_THEN5_I_I]] ]
+// APPROX-NEXT: [[__TAGP_ADDR_0_I28_I_I10:%.*]] = phi ptr [ [[INCDEC_PTR_I34_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ [[INCDEC_PTR_I_I]], %[[IF_THEN5_I_I]] ]
+// APPROX-NEXT: [[TMP4:%.*]] = add i8 [[TMP3]], -48
+// APPROX-NEXT: [[OR_COND_I32_I_I:%.*]] = icmp ult i8 [[TMP4]], 10
+// APPROX-NEXT: br i1 [[OR_COND_I32_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE_I_I_I:.*]]
// APPROX: [[IF_ELSE_I_I_I]]:
-// APPROX-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
-// APPROX-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// APPROX-NEXT: [[TMP5:%.*]] = add i8 [[TMP3]], -97
+// APPROX-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
// APPROX-NEXT: br i1 [[OR_COND33_I_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE17_I_I_I:.*]]
// APPROX: [[IF_ELSE17_I_I_I]]:
-// APPROX-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
-// APPROX-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// APPROX-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[CLEANUP_I36_I_I]]
+// APPROX-NEXT: [[TMP6:%.*]] = add i8 [[TMP3]], -65
+// APPROX-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
+// APPROX-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[_ZL4NANFPKC_EXIT]]
// APPROX: [[IF_END31_I_I_I]]:
-// APPROX-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I34_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
-// APPROX-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4
-// APPROX-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// APPROX-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
+// APPROX-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I29_I_I11]], 4
+// APPROX-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// APPROX-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
-// APPROX-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
-// APPROX-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I31_I_I]], i64 1
-// APPROX-NEXT: br label %[[CLEANUP_I36_I_I]]
-// APPROX: [[CLEANUP_I36_I_I]]:
-// APPROX-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], %[[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// APPROX-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// APPROX-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, %[[IF_END31_I_I_I]] ], [ false, %[[IF_ELSE17_I_I_I]] ]
-// APPROX-NEXT: br i1 [[COND_I_I_I]], label %[[WHILE_COND_I30_I_I]], label %[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP12]]
-// APPROX: [[WHILE_COND_I_I_I]]:
-// APPROX-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], %[[CLEANUP_I_I_I:.*]] ], [ [[INCDEC_PTR_I_I]], %[[IF_THEN_I_I]] ]
-// APPROX-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], %[[CLEANUP_I_I_I]] ], [ 0, %[[IF_THEN_I_I]] ]
-// APPROX-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// APPROX-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
-// APPROX-NEXT: br i1 [[CMP_NOT_I_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I_I_I:.*]]
+// APPROX-NEXT: [[ADD28_I_I_I]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I34_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I28_I_I10]], i64 1
+// APPROX-NEXT: [[TMP7]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// APPROX-NEXT: [[CMP_NOT_I30_I_I:%.*]] = icmp eq i8 [[TMP7]], 0
+// APPROX-NEXT: br i1 [[CMP_NOT_I30_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I31_I_I]], !llvm.loop [[LOOP12]]
// APPROX: [[WHILE_BODY_I_I_I]]:
-// APPROX-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// APPROX-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
-// APPROX-NEXT: br i1 [[OR_COND_I_I_I]], label %[[IF_THEN_I_I_I:.*]], label %[[CLEANUP_I_I_I]]
+// APPROX-NEXT: [[TMP8:%.*]] = phi i8 [ [[TMP10:%.*]], %[[IF_THEN_I_I_I:.*]] ], [ [[TMP2]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[__R_0_I_I_I16:%.*]] = phi i64 [ [[SUB_I_I_I:%.*]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[__TAGP_ADDR_0_I_I_I15:%.*]] = phi ptr [ [[INCDEC_PTR_I_I_I:%.*]], %[[IF_THEN_I_I_I]] ], [ [[INCDEC_PTR_I_I]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[TMP9:%.*]] = and i8 [[TMP8]], -8
+// APPROX-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP9]], 48
+// APPROX-NEXT: br i1 [[OR_COND_I_I_I]], label %[[IF_THEN_I_I_I]], label %[[_ZL4NANFPKC_EXIT]]
// APPROX: [[IF_THEN_I_I_I]]:
-// APPROX-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
-// APPROX-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// APPROX-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I16]], 3
+// APPROX-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
// APPROX-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
-// APPROX-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
-// APPROX-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1
-// APPROX-NEXT: br label %[[CLEANUP_I_I_I]]
-// APPROX: [[CLEANUP_I_I_I]]:
-// APPROX-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], %[[WHILE_BODY_I_I_I]] ]
-// APPROX-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_BODY_I_I_I]] ]
-// APPROX-NEXT: br i1 [[OR_COND_I_I_I]], label %[[WHILE_COND_I_I_I]], label %[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP8]]
-// APPROX: [[WHILE_COND_I14_I_I]]:
-// APPROX-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I_I:%.*]], %[[CLEANUP_I20_I_I:.*]] ], [ [[TAG]], %[[ENTRY]] ]
-// APPROX-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]], %[[CLEANUP_I20_I_I]] ], [ 0, %[[ENTRY]] ]
-// APPROX-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// APPROX-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// APPROX-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I18_I_I:.*]]
+// APPROX-NEXT: [[SUB_I_I_I]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I_I_I15]], i64 1
+// APPROX-NEXT: [[TMP10]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// APPROX-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP10]], 0
+// APPROX-NEXT: br i1 [[CMP_NOT_I_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I_I_I]], !llvm.loop [[LOOP8]]
// APPROX: [[WHILE_BODY_I18_I_I]]:
-// APPROX-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// APPROX-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// APPROX-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[IF_THEN_I24_I_I:.*]], label %[[CLEANUP_I20_I_I]]
-// APPROX: [[IF_THEN_I24_I_I]]:
-// APPROX-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
-// APPROX-NEXT: [[CONV5_I26_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
-// APPROX-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48
-// APPROX-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]], [[CONV5_I26_I_I]]
-// APPROX-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I_I]], i64 1
-// APPROX-NEXT: br label %[[CLEANUP_I20_I_I]]
-// APPROX: [[CLEANUP_I20_I_I]]:
-// APPROX-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [ [[INCDEC_PTR_I29_I_I]], %[[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]], %[[WHILE_BODY_I18_I_I]] ]
-// APPROX-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]], %[[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_BODY_I18_I_I]] ]
-// APPROX-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[WHILE_COND_I14_I_I]], label %[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP11]]
+// APPROX-NEXT: [[TMP11:%.*]] = phi i8 [ [[TMP13:%.*]], %[[IF_THEN_I21_I_I:.*]] ], [ [[TMP1]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[__R_0_I16_I_I7:%.*]] = phi i64 [ [[SUB_I25_I_I:%.*]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[__TAGP_ADDR_0_I15_I_I6:%.*]] = phi ptr [ [[INCDEC_PTR_I26_I_I:%.*]], %[[IF_THEN_I21_I_I]] ], [ [[TAG]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[TMP12:%.*]] = add i8 [[TMP11]], -48
+// APPROX-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP12]], 10
+// APPROX-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[IF_THEN_I21_I_I]], label %[[_ZL4NANFPKC_EXIT]]
+// APPROX: [[IF_THEN_I21_I_I]]:
+// APPROX-NEXT: [[MUL_I22_I_I:%.*]] = mul i64 [[__R_0_I16_I_I7]], 10
+// APPROX-NEXT: [[CONV5_I23_I_I:%.*]] = zext nneg i8 [[TMP11]] to i64
+// APPROX-NEXT: [[ADD_I24_I_I:%.*]] = add i64 [[MUL_I22_I_I]], -48
+// APPROX-NEXT: [[SUB_I25_I_I]] = add i64 [[ADD_I24_I_I]], [[CONV5_I23_I_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I26_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I_I6]], i64 1
+// APPROX-NEXT: [[TMP13]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// APPROX-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP13]], 0
+// APPROX-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I18_I_I]], !llvm.loop [[LOOP11]]
// APPROX: [[_ZL4NANFPKC_EXIT]]:
-// APPROX-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ [[__R_0_I32_I_I]], %[[WHILE_COND_I30_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_COND_I_I_I]] ], [ 0, %[[CLEANUP_I_I_I]] ], [ 0, %[[CLEANUP_I36_I_I]] ], [ 0, %[[CLEANUP_I20_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_COND_I14_I_I]] ]
+// APPROX-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
// APPROX-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32
// APPROX-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303
// APPROX-NEXT: [[BF_SET9_I:%.*]] = or disjoint i32 [[BF_VALUE_I]], 2143289344
-// APPROX-NEXT: [[TMP10:%.*]] = bitcast i32 [[BF_SET9_I]] to float
-// APPROX-NEXT: ret float [[TMP10]]
+// APPROX-NEXT: [[TMP14:%.*]] = bitcast i32 [[BF_SET9_I]] to float
+// APPROX-NEXT: ret float [[TMP14]]
//
// NCRDIV-LABEL: define dso_local float @test_nanf(
// NCRDIV-SAME: ptr noundef readonly captures(none) [[TAG:%.*]]) local_unnamed_addr #[[ATTR2]] {
-// NCRDIV-NEXT: [[ENTRY:.*]]:
+// NCRDIV-NEXT: [[ENTRY:.*:]]
// NCRDIV-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG]], align 1, !tbaa [[CHAR_TBAA7]]
// NCRDIV-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// NCRDIV-NEXT: br i1 [[CMP_I_I]], label %[[IF_THEN_I_I:.*]], label %[[WHILE_COND_I14_I_I:.*]]
+// NCRDIV-NEXT: br i1 [[CMP_I_I]], label %[[IF_THEN_I_I:.*]], label %[[WHILE_COND_I14_I_I_PREHEADER:.*]]
+// NCRDIV: [[WHILE_COND_I14_I_I_PREHEADER]]:
+// NCRDIV-NEXT: [[TMP1:%.*]] = load i8, ptr [[TAG]], align 1, !tbaa [[CHAR_TBAA7]]
+// NCRDIV-NEXT: [[CMP_NOT_I17_I_I5:%.*]] = icmp eq i8 [[TMP1]], 0
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I17_I_I5]], label %[[_ZL4NANFPKC_EXIT:.*]], label %[[WHILE_BODY_I18_I_I:.*]]
// NCRDIV: [[IF_THEN_I_I]]:
// NCRDIV-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[TAG]], i64 1
-// NCRDIV-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// NCRDIV-NEXT: switch i8 [[TMP1]], label %[[WHILE_COND_I_I_I:.*]] [
-// NCRDIV-NEXT: i8 120, label %[[WHILE_COND_I30_I_I_PREHEADER:.*]]
-// NCRDIV-NEXT: i8 88, label %[[WHILE_COND_I30_I_I_PREHEADER]]
+// NCRDIV-NEXT: [[TMP2:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// NCRDIV-NEXT: [[CMP_NOT_I_I_I14:%.*]] = icmp eq i8 [[TMP2]], 0
+// NCRDIV-NEXT: switch i8 [[TMP2]], label %[[WHILE_COND_I_I_I_PREHEADER:.*]] [
+// NCRDIV-NEXT: i8 120, label %[[IF_THEN5_I_I:.*]]
+// NCRDIV-NEXT: i8 88, label %[[IF_THEN5_I_I]]
// NCRDIV-NEXT: ]
-// NCRDIV: [[WHILE_COND_I30_I_I_PREHEADER]]:
-// NCRDIV-NEXT: br label %[[WHILE_COND_I30_I_I:.*]]
-// NCRDIV: [[WHILE_COND_I30_I_I]]:
-// NCRDIV-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I_I:%.*]], %[[CLEANUP_I36_I_I:.*]] ], [ [[INCDEC_PTR_I_I]], %[[WHILE_COND_I30_I_I_PREHEADER]] ]
-// NCRDIV-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], %[[CLEANUP_I36_I_I]] ], [ 0, %[[WHILE_COND_I30_I_I_PREHEADER]] ]
-// NCRDIV-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// NCRDIV-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// NCRDIV-NEXT: br i1 [[CMP_NOT_I33_I_I]], label %[[_ZL4NANFPKC_EXIT:.*]], label %[[WHILE_BODY_I34_I_I:.*]]
-// NCRDIV: [[WHILE_BODY_I34_I_I]]:
-// NCRDIV-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// NCRDIV-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// NCRDIV-NEXT: br i1 [[OR_COND_I35_I_I]], label %[[IF_END31_I_I_I:.*]], label %[[IF_ELSE_I_I_I:.*]]
+// NCRDIV: [[WHILE_COND_I_I_I_PREHEADER]]:
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I_I_I14]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I_I_I:.*]]
+// NCRDIV: [[IF_THEN5_I_I]]:
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I_I_I14]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I31_I_I:.*]]
+// NCRDIV: [[WHILE_BODY_I31_I_I]]:
+// NCRDIV-NEXT: [[TMP3:%.*]] = phi i8 [ [[TMP7:%.*]], %[[IF_END31_I_I_I:.*]] ], [ [[TMP2]], %[[IF_THEN5_I_I]] ]
+// NCRDIV-NEXT: [[__R_0_I29_I_I11:%.*]] = phi i64 [ [[ADD28_I_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ 0, %[[IF_THEN5_I_I]] ]
+// NCRDIV-NEXT: [[__TAGP_ADDR_0_I28_I_I10:%.*]] = phi ptr [ [[INCDEC_PTR_I34_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ [[INCDEC_PTR_I_I]], %[[IF_THEN5_I_I]] ]
+// NCRDIV-NEXT: [[TMP4:%.*]] = add i8 [[TMP3]], -48
+// NCRDIV-NEXT: [[OR_COND_I32_I_I:%.*]] = icmp ult i8 [[TMP4]], 10
+// NCRDIV-NEXT: br i1 [[OR_COND_I32_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE_I_I_I:.*]]
// NCRDIV: [[IF_ELSE_I_I_I]]:
-// NCRDIV-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
-// NCRDIV-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// NCRDIV-NEXT: [[TMP5:%.*]] = add i8 [[TMP3]], -97
+// NCRDIV-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
// NCRDIV-NEXT: br i1 [[OR_COND33_I_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE17_I_I_I:.*]]
// NCRDIV: [[IF_ELSE17_I_I_I]]:
-// NCRDIV-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
-// NCRDIV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// NCRDIV-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[CLEANUP_I36_I_I]]
+// NCRDIV-NEXT: [[TMP6:%.*]] = add i8 [[TMP3]], -65
+// NCRDIV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
+// NCRDIV-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[_ZL4NANFPKC_EXIT]]
// NCRDIV: [[IF_END31_I_I_I]]:
-// NCRDIV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I34_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
-// NCRDIV-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4
-// NCRDIV-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// NCRDIV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
+// NCRDIV-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I29_I_I11]], 4
+// NCRDIV-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// NCRDIV-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
-// NCRDIV-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
-// NCRDIV-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I31_I_I]], i64 1
-// NCRDIV-NEXT: br label %[[CLEANUP_I36_I_I]]
-// NCRDIV: [[CLEANUP_I36_I_I]]:
-// NCRDIV-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], %[[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// NCRDIV-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// NCRDIV-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, %[[IF_END31_I_I_I]] ], [ false, %[[IF_ELSE17_I_I_I]] ]
-// NCRDIV-NEXT: br i1 [[COND_I_I_I]], label %[[WHILE_COND_I30_I_I]], label %[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP12]]
-// NCRDIV: [[WHILE_COND_I_I_I]]:
-// NCRDIV-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]] ]
-// NCRDIV-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], %[[CLEANUP_I_I_I]] ], [ 0, %[[IF_THEN_I_I]] ]
-// NCRDIV-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// NCRDIV-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
-// NCRDIV-NEXT: br i1 [[CMP_NOT_I_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I_I_I:.*]]
+// NCRDIV-NEXT: [[ADD28_I_I_I]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I34_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I28_I_I10]], i64 1
+// NCRDIV-NEXT: [[TMP7]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// NCRDIV-NEXT: [[CMP_NOT_I30_I_I:%.*]] = icmp eq i8 [[TMP7]], 0
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I30_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I31_I_I]], !llvm.loop [[LOOP12]]
// NCRDIV: [[WHILE_BODY_I_I_I]]:
-// NCRDIV-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// NCRDIV-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
-// NCRDIV-NEXT: br i1 [[OR_COND_I_I_I]], label %[[IF_THEN_I_I_I:.*]], label %[[CLEANUP_I_I_I]]
+// NCRDIV-NEXT: [[TMP8:%.*]] = phi i8 [ [[TMP10:%.*]], %[[IF_THEN_I_I_I:.*]] ], [ [[TMP2]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[__R_0_I_I_I16:%.*]] = phi i64 [ [[SUB_I_I_I:%.*]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[__TAGP_ADDR_0_I_I_I15:%.*]] = phi ptr [ [[INCDEC_PTR_I_I_I:%.*]], %[[IF_THEN_I_I_I]] ], [ [[INCDEC_PTR_I_I]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[TMP9:%.*]] = and i8 [[TMP8]], -8
+// NCRDIV-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP9]], 48
+// NCRDIV-NEXT: br i1 [[OR_COND_I_I_I]], label %[[IF_THEN_I_I_I]], label %[[_ZL4NANFPKC_EXIT]]
// NCRDIV: [[IF_THEN_I_I_I]]:
-// NCRDIV-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
-// NCRDIV-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// NCRDIV-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I16]], 3
+// NCRDIV-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
// NCRDIV-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
-// NCRDIV-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
-// NCRDIV-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1
-// NCRDIV-NEXT: br label %[[CLEANUP_I_I_I]]
-// NCRDIV: [[CLEANUP_I_I_I]]:
-// NCRDIV-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]] ]
-// NCRDIV-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]] ]
-// NCRDIV-NEXT: br i1 [[OR_COND_I_I_I]], label %[[WHILE_COND_I_I_I]], label %[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP8]]
-// NCRDIV: [[WHILE_COND_I14_I_I]]:
-// NCRDIV-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I_I:%.*]], %[[CLEANUP_I20_I_I:.*]] ], [ [[TAG]], %[[ENTRY]] ]
-// NCRDIV-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]], %[[CLEANUP_I20_I_I]] ], [ 0, %[[ENTRY]] ]
-// NCRDIV-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// NCRDIV-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// NCRDIV-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I18_I_I:.*]]
+// NCRDIV-NEXT: [[SUB_I_I_I]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I_I_I15]], i64 1
+// NCRDIV-NEXT: [[TMP10]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// NCRDIV-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP10]], 0
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I_I_I]], !llvm.loop [[LOOP8]]
// NCRDIV: [[WHILE_BODY_I18_I_I]]:
-// NCRDIV-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// NCRDIV-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// NCRDIV-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[IF_THEN_I24_I_I:.*]], label %[[CLEANUP_I20_I_I]]
-// NCRDIV: [[IF_THEN_I24_I_I]]:
-// NCRDIV-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
-// NCRDIV-NEXT: [[CONV5_I26_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
-// NCRDIV-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48
-// NCRDIV-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]], [[CONV5_I26_I_I]]
-// NCRDIV-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I_I]], i64 1
-// NCRDIV-NEXT: br label %[[CLEANUP_I20_I_I]]
-// NCRDIV: [[CLEANUP_I20_I_I]]:
-// NCRDIV-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [ [[INCDEC_PTR_I29_I_I]], %[[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]], %[[WHILE_BODY_I18_I_I]] ]
-// NCRDIV-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]], %[[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_BODY_I18_I_I]] ]
-// NCRDIV-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[WHILE_COND_I14_I_I]], label %[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP11]]
+// NCRDIV-NEXT: [[TMP11:%.*]] = phi i8 [ [[TMP13:%.*]], %[[IF_THEN_I21_I_I:.*]] ], [ [[TMP1]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[__R_0_I16_I_I7:%.*]] = phi i64 [ [[SUB_I25_I_I:%.*]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[__TAGP_ADDR_0_I15_I_I6:%.*]] = phi ptr [ [[INCDEC_PTR_I26_I_I:%.*]], %[[IF_THEN_I21_I_I]] ], [ [[TAG]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[TMP12:%.*]] = add i8 [[TMP11]], -48
+// NCRDIV-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP12]], 10
+// NCRDIV-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[IF_THEN_I21_I_I]], label %[[_ZL4NANFPKC_EXIT]]
+// NCRDIV: [[IF_THEN_I21_I_I]]:
+// NCRDIV-NEXT: [[MUL_I22_I_I:%.*]] = mul i64 [[__R_0_I16_I_I7]], 10
+// NCRDIV-NEXT: [[CONV5_I23_I_I:%.*]] = zext nneg i8 [[TMP11]] to i64
+// NCRDIV-NEXT: [[ADD_I24_I_I:%.*]] = add i64 [[MUL_I22_I_I]], -48
+// NCRDIV-NEXT: [[SUB_I25_I_I]] = add i64 [[ADD_I24_I_I]], [[CONV5_I23_I_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I26_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I_I6]], i64 1
+// NCRDIV-NEXT: [[TMP13]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// NCRDIV-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP13]], 0
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I18_I_I]], !llvm.loop [[LOOP11]]
// NCRDIV: [[_ZL4NANFPKC_EXIT]]:
-// NCRDIV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ [[__R_0_I32_I_I]], %[[WHILE_COND_I30_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_COND_I_I_I]] ], [ 0, %[[CLEANUP_I_I_I]] ], [ 0, %[[CLEANUP_I36_I_I]] ], [ 0, %[[CLEANUP_I20_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_COND_I14_I_I]] ]
+// NCRDIV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
// NCRDIV-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32
// NCRDIV-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303
// NCRDIV-NEXT: [[BF_SET9_I:%.*]] = or disjoint i32 [[BF_VALUE_I]], 2143289344
-// NCRDIV-NEXT: [[TMP10:%.*]] = bitcast i32 [[BF_SET9_I]] to float
-// NCRDIV-NEXT: ret float [[TMP10]]
+// NCRDIV-NEXT: [[TMP14:%.*]] = bitcast i32 [[BF_SET9_I]] to float
+// NCRDIV-NEXT: ret float [[TMP14]]
//
// AMDGCNSPIRV-LABEL: define spir_func float @test_nanf(
// AMDGCNSPIRV-SAME: ptr addrspace(4) noundef readonly captures(none) [[TAG:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR2]] {
@@ -4650,53 +4596,49 @@ extern "C" __device__ double test_modf(double x, double* y) {
// AMDGCNSPIRV-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[TAG]], i64 1
// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = load i8, ptr addrspace(4) [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
// AMDGCNSPIRV-NEXT: switch i8 [[TMP1]], label %[[WHILE_COND_I_I_I:.*]] [
-// AMDGCNSPIRV-NEXT: i8 120, label %[[WHILE_COND_I28_I_I_PREHEADER:.*]]
-// AMDGCNSPIRV-NEXT: i8 88, label %[[WHILE_COND_I28_I_I_PREHEADER]]
+// AMDGCNSPIRV-NEXT: i8 120, label %[[IF_THEN5_I_I:.*]]
+// AMDGCNSPIRV-NEXT: i8 88, label %[[IF_THEN5_I_I]]
// AMDGCNSPIRV-NEXT: ]
-// AMDGCNSPIRV: [[WHILE_COND_I28_I_I_PREHEADER]]:
-// AMDGCNSPIRV-NEXT: br label %[[WHILE_COND_I28_I_I:.*]]
-// AMDGCNSPIRV: [[WHILE_COND_I28_I_I]]:
-// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_0_I29_I_I:%.*]] = phi ptr addrspace(4) [ [[__TAGP_ADDR_1_I34_I_I:%.*]], %[[CLEANUP_I_I_I:.*]] ], [ [[INCDEC_PTR_I_I]], %[[WHILE_COND_I28_I_I_PREHEADER]] ]
-// AMDGCNSPIRV-NEXT: [[__R_0_I30_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], %[[CLEANUP_I_I_I]] ], [ 0, %[[WHILE_COND_I28_I_I_PREHEADER]] ]
-// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I29_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// AMDGCNSPIRV-NEXT: [[CMP_NOT_I31_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I31_I_I]], label %[[_ZL4NANFPKC_EXIT:.*]], label %[[WHILE_BODY_I32_I_I:.*]]
+// AMDGCNSPIRV: [[IF_THEN5_I_I]]:
+// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = load i8, ptr addrspace(4) [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// AMDGCNSPIRV-NEXT: [[CMP_NOT_I31_I_I5:%.*]] = icmp eq i8 [[TMP2]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I31_I_I5]], label %[[_ZL4NANFPKC_EXIT:.*]], label %[[WHILE_BODY_I32_I_I:.*]]
// AMDGCNSPIRV: [[WHILE_BODY_I32_I_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// AMDGCNSPIRV-NEXT: [[OR_COND_I33_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// AMDGCNSPIRV-NEXT: br i1 [[OR_COND_I33_I_I]], label %[[IF_END31_I_I_I:.*]], label %[[IF_ELSE_I_I_I:.*]]
+// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = phi i8 [ [[TMP7:%.*]], %[[IF_END31_I_I_I:.*]] ], [ [[TMP2]], %[[IF_THEN5_I_I]] ]
+// AMDGCNSPIRV-NEXT: [[__R_0_I30_I_I7:%.*]] = phi i64 [ [[ADD28_I_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ 0, %[[IF_THEN5_I_I]] ]
+// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_0_I29_I_I6:%.*]] = phi ptr addrspace(4) [ [[INCDEC_PTR_I36_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ [[INCDEC_PTR_I_I]], %[[IF_THEN5_I_I]] ]
+// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = add i8 [[TMP3]], -48
+// AMDGCNSPIRV-NEXT: [[OR_COND_I33_I_I:%.*]] = icmp ult i8 [[TMP4]], 10
+// AMDGCNSPIRV-NEXT: br i1 [[OR_COND_I33_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE_I_I_I:.*]]
// AMDGCNSPIRV: [[IF_ELSE_I_I_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
-// AMDGCNSPIRV-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = add i8 [[TMP3]], -97
+// AMDGCNSPIRV-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
// AMDGCNSPIRV-NEXT: br i1 [[OR_COND33_I_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE17_I_I_I:.*]]
// AMDGCNSPIRV: [[IF_ELSE17_I_I_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
-// AMDGCNSPIRV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// AMDGCNSPIRV-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[CLEANUP_I_I_I]]
+// AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = add i8 [[TMP3]], -65
+// AMDGCNSPIRV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
+// AMDGCNSPIRV-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[_ZL4NANFPKC_EXIT]]
// AMDGCNSPIRV: [[IF_END31_I_I_I]]:
// AMDGCNSPIRV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I32_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
-// AMDGCNSPIRV-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I30_I_I]], 4
-// AMDGCNSPIRV-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// AMDGCNSPIRV-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I30_I_I7]], 4
+// AMDGCNSPIRV-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// AMDGCNSPIRV-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
-// AMDGCNSPIRV-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
-// AMDGCNSPIRV-NEXT: [[INCDEC_PTR_I37_I_I:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__TAGP_ADDR_0_I29_I_I]], i64 1
-// AMDGCNSPIRV-NEXT: br label %[[CLEANUP_I_I_I]]
-// AMDGCNSPIRV: [[CLEANUP_I_I_I]]:
-// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_1_I34_I_I]] = phi ptr addrspace(4) [ [[INCDEC_PTR_I37_I_I]], %[[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I29_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// AMDGCNSPIRV-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ [[__R_0_I30_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// AMDGCNSPIRV-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, %[[IF_END31_I_I_I]] ], [ false, %[[IF_ELSE17_I_I_I]] ]
-// AMDGCNSPIRV-NEXT: br i1 [[COND_I_I_I]], label %[[WHILE_COND_I28_I_I]], label %[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP13]]
+// AMDGCNSPIRV-NEXT: [[ADD28_I_I_I]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
+// AMDGCNSPIRV-NEXT: [[INCDEC_PTR_I36_I_I]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__TAGP_ADDR_0_I29_I_I6]], i64 1
+// AMDGCNSPIRV-NEXT: [[TMP7]] = load i8, ptr addrspace(4) [[INCDEC_PTR_I36_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// AMDGCNSPIRV-NEXT: [[CMP_NOT_I31_I_I:%.*]] = icmp eq i8 [[TMP7]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I31_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I32_I_I]], !llvm.loop [[LOOP13]]
// AMDGCNSPIRV: [[WHILE_COND_I_I_I]]:
// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr addrspace(4) [ [[__TAGP_ADDR_1_I_I_I:%.*]], %[[WHILE_BODY_I_I_I:.*]] ], [ [[INCDEC_PTR_I_I]], %[[IF_THEN_I_I]] ]
// AMDGCNSPIRV-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_THEN_I_I]] ]
-// AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// AMDGCNSPIRV-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// AMDGCNSPIRV-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I_I_I]]
// AMDGCNSPIRV: [[WHILE_BODY_I_I_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// AMDGCNSPIRV-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = and i8 [[TMP8]], -8
+// AMDGCNSPIRV-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP9]], 48
// AMDGCNSPIRV-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
-// AMDGCNSPIRV-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// AMDGCNSPIRV-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
// AMDGCNSPIRV-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
// AMDGCNSPIRV-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_1_I_I_I_IDX:%.*]] = zext i1 [[OR_COND_I_I_I]] to i64
@@ -4706,14 +4648,14 @@ extern "C" __device__ double test_modf(double x, double* y) {
// AMDGCNSPIRV: [[WHILE_COND_I14_I_I]]:
// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr addrspace(4) [ [[__TAGP_ADDR_1_I25_I_I:%.*]], %[[WHILE_BODY_I18_I_I:.*]] ], [ [[TAG]], %[[ENTRY]] ]
// AMDGCNSPIRV-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I26_I_I:%.*]], %[[WHILE_BODY_I18_I_I]] ], [ 0, %[[ENTRY]] ]
-// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// AMDGCNSPIRV-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
+// AMDGCNSPIRV-NEXT: [[TMP10:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// AMDGCNSPIRV-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP10]], 0
// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I18_I_I]]
// AMDGCNSPIRV: [[WHILE_BODY_I18_I_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// AMDGCNSPIRV-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
+// AMDGCNSPIRV-NEXT: [[TMP11:%.*]] = add i8 [[TMP10]], -48
+// AMDGCNSPIRV-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP11]], 10
// AMDGCNSPIRV-NEXT: [[MUL_I20_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
-// AMDGCNSPIRV-NEXT: [[CONV5_I21_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
+// AMDGCNSPIRV-NEXT: [[CONV5_I21_I_I:%.*]] = zext nneg i8 [[TMP10]] to i64
// AMDGCNSPIRV-NEXT: [[ADD_I22_I_I:%.*]] = add i64 [[MUL_I20_I_I]], -48
// AMDGCNSPIRV-NEXT: [[SUB_I23_I_I:%.*]] = add i64 [[ADD_I22_I_I]], [[CONV5_I21_I_I]]
// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_1_I25_I_I_IDX:%.*]] = zext i1 [[OR_COND_I19_I_I]] to i64
@@ -4721,12 +4663,12 @@ extern "C" __device__ double test_modf(double x, double* y) {
// AMDGCNSPIRV-NEXT: [[__R_1_I26_I_I]] = select i1 [[OR_COND_I19_I_I]], i64 [[SUB_I23_I_I]], i64 [[__R_0_I16_I_I]]
// AMDGCNSPIRV-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[WHILE_COND_I14_I_I]], label %[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP12]]
// AMDGCNSPIRV: [[_ZL4NANFPKC_EXIT]]:
-// AMDGCNSPIRV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ [[__R_0_I30_I_I]], %[[WHILE_COND_I28_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_COND_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[CLEANUP_I_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_COND_I14_I_I]] ]
+// AMDGCNSPIRV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_COND_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_COND_I14_I_I]] ]
// AMDGCNSPIRV-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32
// AMDGCNSPIRV-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303
// AMDGCNSPIRV-NEXT: [[BF_SET9_I:%.*]] = or disjoint i32 [[BF_VALUE_I]], 2143289344
-// AMDGCNSPIRV-NEXT: [[TMP10:%.*]] = bitcast i32 [[BF_SET9_I]] to float
-// AMDGCNSPIRV-NEXT: ret float [[TMP10]]
+// AMDGCNSPIRV-NEXT: [[TMP12:%.*]] = bitcast i32 [[BF_SET9_I]] to float
+// AMDGCNSPIRV-NEXT: ret float [[TMP12]]
//
extern "C" __device__ float test_nanf(const char *tag) {
return nanf(tag);
@@ -4734,98 +4676,89 @@ extern "C" __device__ float test_nanf(const char *tag) {
// DEFAULT-LABEL: define dso_local double @test_nan(
// DEFAULT-SAME: ptr noundef readonly captures(none) [[TAG:%.*]]) local_unnamed_addr #[[ATTR2]] {
-// DEFAULT-NEXT: [[ENTRY:.*]]:
+// DEFAULT-NEXT: [[ENTRY:.*:]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG]], align 1, !tbaa [[CHAR_TBAA7]]
// DEFAULT-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// DEFAULT-NEXT: br i1 [[CMP_I_I]], label %[[IF_THEN_I_I:.*]], label %[[WHILE_COND_I14_I_I:.*]]
+// DEFAULT-NEXT: br i1 [[CMP_I_I]], label %[[IF_THEN_I_I:.*]], label %[[WHILE_COND_I14_I_I_PREHEADER:.*]]
+// DEFAULT: [[WHILE_COND_I14_I_I_PREHEADER]]:
+// DEFAULT-NEXT: [[TMP1:%.*]] = load i8, ptr [[TAG]], align 1, !tbaa [[CHAR_TBAA7]]
+// DEFAULT-NEXT: [[CMP_NOT_I17_I_I5:%.*]] = icmp eq i8 [[TMP1]], 0
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I17_I_I5]], label %[[_ZL3NANPKC_EXIT:.*]], label %[[WHILE_BODY_I18_I_I:.*]]
// DEFAULT: [[IF_THEN_I_I]]:
// DEFAULT-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[TAG]], i64 1
-// DEFAULT-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// DEFAULT-NEXT: switch i8 [[TMP1]], label %[[WHILE_COND_I_I_I:.*]] [
-// DEFAULT-NEXT: i8 120, label %[[WHILE_COND_I30_I_I_PREHEADER:.*]]
-// DEFAULT-NEXT: i8 88, label %[[WHILE_COND_I30_I_I_PREHEADER]]
+// DEFAULT-NEXT: [[TMP2:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// DEFAULT-NEXT: [[CMP_NOT_I_I_I14:%.*]] = icmp eq i8 [[TMP2]], 0
+// DEFAULT-NEXT: switch i8 [[TMP2]], label %[[WHILE_COND_I_I_I_PREHEADER:.*]] [
+// DEFAULT-NEXT: i8 120, label %[[IF_THEN5_I_I:.*]]
+// DEFAULT-NEXT: i8 88, label %[[IF_THEN5_I_I]]
// DEFAULT-NEXT: ]
-// DEFAULT: [[WHILE_COND_I30_I_I_PREHEADER]]:
-// DEFAULT-NEXT: br label %[[WHILE_COND_I30_I_I:.*]]
-// DEFAULT: [[WHILE_COND_I30_I_I]]:
-// DEFAULT-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I_I:%.*]], %[[CLEANUP_I36_I_I:.*]] ], [ [[INCDEC_PTR_I_I]], %[[WHILE_COND_I30_I_I_PREHEADER]] ]
-// DEFAULT-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], %[[CLEANUP_I36_I_I]] ], [ 0, %[[WHILE_COND_I30_I_I_PREHEADER]] ]
-// DEFAULT-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// DEFAULT-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// DEFAULT-NEXT: br i1 [[CMP_NOT_I33_I_I]], label %[[_ZL3NANPKC_EXIT:.*]], label %[[WHILE_BODY_I34_I_I:.*]]
-// DEFAULT: [[WHILE_BODY_I34_I_I]]:
-// DEFAULT-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// DEFAULT-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// DEFAULT-NEXT: br i1 [[OR_COND_I35_I_I]], label %[[IF_END31_I_I_I:.*]], label %[[IF_ELSE_I_I_I:.*]]
+// DEFAULT: [[WHILE_COND_I_I_I_PREHEADER]]:
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I_I_I14]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I_I_I:.*]]
+// DEFAULT: [[IF_THEN5_I_I]]:
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I_I_I14]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I31_I_I:.*]]
+// DEFAULT: [[WHILE_BODY_I31_I_I]]:
+// DEFAULT-NEXT: [[TMP3:%.*]] = phi i8 [ [[TMP7:%.*]], %[[IF_END31_I_I_I:.*]] ], [ [[TMP2]], %[[IF_THEN5_I_I]] ]
+// DEFAULT-NEXT: [[__R_0_I29_I_I11:%.*]] = phi i64 [ [[ADD28_I_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ 0, %[[IF_THEN5_I_I]] ]
+// DEFAULT-NEXT: [[__TAGP_ADDR_0_I28_I_I10:%.*]] = phi ptr [ [[INCDEC_PTR_I34_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ [[INCDEC_PTR_I_I]], %[[IF_THEN5_I_I]] ]
+// DEFAULT-NEXT: [[TMP4:%.*]] = add i8 [[TMP3]], -48
+// DEFAULT-NEXT: [[OR_COND_I32_I_I:%.*]] = icmp ult i8 [[TMP4]], 10
+// DEFAULT-NEXT: br i1 [[OR_COND_I32_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE_I_I_I:.*]]
// DEFAULT: [[IF_ELSE_I_I_I]]:
-// DEFAULT-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
-// DEFAULT-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// DEFAULT-NEXT: [[TMP5:%.*]] = add i8 [[TMP3]], -97
+// DEFAULT-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
// DEFAULT-NEXT: br i1 [[OR_COND33_I_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE17_I_I_I:.*]]
// DEFAULT: [[IF_ELSE17_I_I_I]]:
-// DEFAULT-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
-// DEFAULT-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// DEFAULT-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[CLEANUP_I36_I_I]]
+// DEFAULT-NEXT: [[TMP6:%.*]] = add i8 [[TMP3]], -65
+// DEFAULT-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
+// DEFAULT-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[_ZL3NANPKC_EXIT]]
// DEFAULT: [[IF_END31_I_I_I]]:
-// DEFAULT-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I34_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
-// DEFAULT-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4
-// DEFAULT-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// DEFAULT-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
+// DEFAULT-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I29_I_I11]], 4
+// DEFAULT-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// DEFAULT-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
-// DEFAULT-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
-// DEFAULT-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I31_I_I]], i64 1
-// DEFAULT-NEXT: br label %[[CLEANUP_I36_I_I]]
-// DEFAULT: [[CLEANUP_I36_I_I]]:
-// DEFAULT-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], %[[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// DEFAULT-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// DEFAULT-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, %[[IF_END31_I_I_I]] ], [ false, %[[IF_ELSE17_I_I_I]] ]
-// DEFAULT-NEXT: br i1 [[COND_I_I_I]], label %[[WHILE_COND_I30_I_I]], label %[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP12]]
-// DEFAULT: [[WHILE_COND_I_I_I]]:
-// DEFAULT-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], %[[CLEANUP_I_I_I:.*]] ], [ [[INCDEC_PTR_I_I]], %[[IF_THEN_I_I]] ]
-// DEFAULT-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], %[[CLEANUP_I_I_I]] ], [ 0, %[[IF_THEN_I_I]] ]
-// DEFAULT-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// DEFAULT-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
-// DEFAULT-NEXT: br i1 [[CMP_NOT_I_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I_I_I:.*]]
+// DEFAULT-NEXT: [[ADD28_I_I_I]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I34_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I28_I_I10]], i64 1
+// DEFAULT-NEXT: [[TMP7]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// DEFAULT-NEXT: [[CMP_NOT_I30_I_I:%.*]] = icmp eq i8 [[TMP7]], 0
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I30_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I31_I_I]], !llvm.loop [[LOOP12]]
// DEFAULT: [[WHILE_BODY_I_I_I]]:
-// DEFAULT-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// DEFAULT-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
-// DEFAULT-NEXT: br i1 [[OR_COND_I_I_I]], label %[[IF_THEN_I_I_I:.*]], label %[[CLEANUP_I_I_I]]
+// DEFAULT-NEXT: [[TMP8:%.*]] = phi i8 [ [[TMP10:%.*]], %[[IF_THEN_I_I_I:.*]] ], [ [[TMP2]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[__R_0_I_I_I16:%.*]] = phi i64 [ [[SUB_I_I_I:%.*]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[__TAGP_ADDR_0_I_I_I15:%.*]] = phi ptr [ [[INCDEC_PTR_I_I_I:%.*]], %[[IF_THEN_I_I_I]] ], [ [[INCDEC_PTR_I_I]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[TMP9:%.*]] = and i8 [[TMP8]], -8
+// DEFAULT-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP9]], 48
+// DEFAULT-NEXT: br i1 [[OR_COND_I_I_I]], label %[[IF_THEN_I_I_I]], label %[[_ZL3NANPKC_EXIT]]
// DEFAULT: [[IF_THEN_I_I_I]]:
-// DEFAULT-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
-// DEFAULT-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// DEFAULT-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I16]], 3
+// DEFAULT-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
// DEFAULT-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
-// DEFAULT-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
-// DEFAULT-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1
-// DEFAULT-NEXT: br label %[[CLEANUP_I_I_I]]
-// DEFAULT: [[CLEANUP_I_I_I]]:
-// DEFAULT-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], %[[WHILE_BODY_I_I_I]] ]
-// DEFAULT-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_BODY_I_I_I]] ]
-// DEFAULT-NEXT: br i1 [[OR_COND_I_I_I]], label %[[WHILE_COND_I_I_I]], label %[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP8]]
-// DEFAULT: [[WHILE_COND_I14_I_I]]:
-// DEFAULT-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I_I:%.*]], %[[CLEANUP_I20_I_I:.*]] ], [ [[TAG]], %[[ENTRY]] ]
-// DEFAULT-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]], %[[CLEANUP_I20_I_I]] ], [ 0, %[[ENTRY]] ]
-// DEFAULT-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// DEFAULT-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// DEFAULT-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I18_I_I:.*]]
+// DEFAULT-NEXT: [[SUB_I_I_I]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I_I_I15]], i64 1
+// DEFAULT-NEXT: [[TMP10]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// DEFAULT-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP10]], 0
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I_I_I]], !llvm.loop [[LOOP8]]
// DEFAULT: [[WHILE_BODY_I18_I_I]]:
-// DEFAULT-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// DEFAULT-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// DEFAULT-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[IF_THEN_I24_I_I:.*]], label %[[CLEANUP_I20_I_I]]
-// DEFAULT: [[IF_THEN_I24_I_I]]:
-// DEFAULT-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
-// DEFAULT-NEXT: [[CONV5_I26_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
-// DEFAULT-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48
-// DEFAULT-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]], [[CONV5_I26_I_I]]
-// DEFAULT-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I_I]], i64 1
-// DEFAULT-NEXT: br label %[[CLEANUP_I20_I_I]]
-// DEFAULT: [[CLEANUP_I20_I_I]]:
-// DEFAULT-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [ [[INCDEC_PTR_I29_I_I]], %[[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]], %[[WHILE_BODY_I18_I_I]] ]
-// DEFAULT-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]], %[[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_BODY_I18_I_I]] ]
-// DEFAULT-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[WHILE_COND_I14_I_I]], label %[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP11]]
+// DEFAULT-NEXT: [[TMP11:%.*]] = phi i8 [ [[TMP13:%.*]], %[[IF_THEN_I21_I_I:.*]] ], [ [[TMP1]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[__R_0_I16_I_I7:%.*]] = phi i64 [ [[SUB_I25_I_I:%.*]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[__TAGP_ADDR_0_I15_I_I6:%.*]] = phi ptr [ [[INCDEC_PTR_I26_I_I:%.*]], %[[IF_THEN_I21_I_I]] ], [ [[TAG]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// DEFAULT-NEXT: [[TMP12:%.*]] = add i8 [[TMP11]], -48
+// DEFAULT-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP12]], 10
+// DEFAULT-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[IF_THEN_I21_I_I]], label %[[_ZL3NANPKC_EXIT]]
+// DEFAULT: [[IF_THEN_I21_I_I]]:
+// DEFAULT-NEXT: [[MUL_I22_I_I:%.*]] = mul i64 [[__R_0_I16_I_I7]], 10
+// DEFAULT-NEXT: [[CONV5_I23_I_I:%.*]] = zext nneg i8 [[TMP11]] to i64
+// DEFAULT-NEXT: [[ADD_I24_I_I:%.*]] = add i64 [[MUL_I22_I_I]], -48
+// DEFAULT-NEXT: [[SUB_I25_I_I]] = add i64 [[ADD_I24_I_I]], [[CONV5_I23_I_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I26_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I_I6]], i64 1
+// DEFAULT-NEXT: [[TMP13]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// DEFAULT-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP13]], 0
+// DEFAULT-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I18_I_I]], !llvm.loop [[LOOP11]]
// DEFAULT: [[_ZL3NANPKC_EXIT]]:
-// DEFAULT-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ [[__R_0_I32_I_I]], %[[WHILE_COND_I30_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_COND_I_I_I]] ], [ 0, %[[CLEANUP_I_I_I]] ], [ 0, %[[CLEANUP_I36_I_I]] ], [ 0, %[[CLEANUP_I20_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_COND_I14_I_I]] ]
+// DEFAULT-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
// DEFAULT-NEXT: [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]], 2251799813685247
// DEFAULT-NEXT: [[BF_SET9_I:%.*]] = or disjoint i64 [[BF_VALUE_I]], 9221120237041090560
-// DEFAULT-NEXT: [[TMP10:%.*]] = bitcast i64 [[BF_SET9_I]] to double
-// DEFAULT-NEXT: ret double [[TMP10]]
+// DEFAULT-NEXT: [[TMP14:%.*]] = bitcast i64 [[BF_SET9_I]] to double
+// DEFAULT-NEXT: ret double [[TMP14]]
//
// FINITEONLY-LABEL: define dso_local nofpclass(nan inf) double @test_nan(
// FINITEONLY-SAME: ptr noundef readonly captures(none) [[TAG:%.*]]) local_unnamed_addr #[[ATTR3]] {
@@ -4834,193 +4767,175 @@ extern "C" __device__ float test_nanf(const char *tag) {
//
// APPROX-LABEL: define dso_local double @test_nan(
// APPROX-SAME: ptr noundef readonly captures(none) [[TAG:%.*]]) local_unnamed_addr #[[ATTR2]] {
-// APPROX-NEXT: [[ENTRY:.*]]:
+// APPROX-NEXT: [[ENTRY:.*:]]
// APPROX-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG]], align 1, !tbaa [[CHAR_TBAA7]]
// APPROX-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// APPROX-NEXT: br i1 [[CMP_I_I]], label %[[IF_THEN_I_I:.*]], label %[[WHILE_COND_I14_I_I:.*]]
+// APPROX-NEXT: br i1 [[CMP_I_I]], label %[[IF_THEN_I_I:.*]], label %[[WHILE_COND_I14_I_I_PREHEADER:.*]]
+// APPROX: [[WHILE_COND_I14_I_I_PREHEADER]]:
+// APPROX-NEXT: [[TMP1:%.*]] = load i8, ptr [[TAG]], align 1, !tbaa [[CHAR_TBAA7]]
+// APPROX-NEXT: [[CMP_NOT_I17_I_I5:%.*]] = icmp eq i8 [[TMP1]], 0
+// APPROX-NEXT: br i1 [[CMP_NOT_I17_I_I5]], label %[[_ZL3NANPKC_EXIT:.*]], label %[[WHILE_BODY_I18_I_I:.*]]
// APPROX: [[IF_THEN_I_I]]:
// APPROX-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[TAG]], i64 1
-// APPROX-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// APPROX-NEXT: switch i8 [[TMP1]], label %[[WHILE_COND_I_I_I:.*]] [
-// APPROX-NEXT: i8 120, label %[[WHILE_COND_I30_I_I_PREHEADER:.*]]
-// APPROX-NEXT: i8 88, label %[[WHILE_COND_I30_I_I_PREHEADER]]
+// APPROX-NEXT: [[TMP2:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// APPROX-NEXT: [[CMP_NOT_I_I_I14:%.*]] = icmp eq i8 [[TMP2]], 0
+// APPROX-NEXT: switch i8 [[TMP2]], label %[[WHILE_COND_I_I_I_PREHEADER:.*]] [
+// APPROX-NEXT: i8 120, label %[[IF_THEN5_I_I:.*]]
+// APPROX-NEXT: i8 88, label %[[IF_THEN5_I_I]]
// APPROX-NEXT: ]
-// APPROX: [[WHILE_COND_I30_I_I_PREHEADER]]:
-// APPROX-NEXT: br label %[[WHILE_COND_I30_I_I:.*]]
-// APPROX: [[WHILE_COND_I30_I_I]]:
-// APPROX-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I_I:%.*]], %[[CLEANUP_I36_I_I:.*]] ], [ [[INCDEC_PTR_I_I]], %[[WHILE_COND_I30_I_I_PREHEADER]] ]
-// APPROX-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], %[[CLEANUP_I36_I_I]] ], [ 0, %[[WHILE_COND_I30_I_I_PREHEADER]] ]
-// APPROX-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// APPROX-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// APPROX-NEXT: br i1 [[CMP_NOT_I33_I_I]], label %[[_ZL3NANPKC_EXIT:.*]], label %[[WHILE_BODY_I34_I_I:.*]]
-// APPROX: [[WHILE_BODY_I34_I_I]]:
-// APPROX-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// APPROX-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// APPROX-NEXT: br i1 [[OR_COND_I35_I_I]], label %[[IF_END31_I_I_I:.*]], label %[[IF_ELSE_I_I_I:.*]]
+// APPROX: [[WHILE_COND_I_I_I_PREHEADER]]:
+// APPROX-NEXT: br i1 [[CMP_NOT_I_I_I14]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I_I_I:.*]]
+// APPROX: [[IF_THEN5_I_I]]:
+// APPROX-NEXT: br i1 [[CMP_NOT_I_I_I14]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I31_I_I:.*]]
+// APPROX: [[WHILE_BODY_I31_I_I]]:
+// APPROX-NEXT: [[TMP3:%.*]] = phi i8 [ [[TMP7:%.*]], %[[IF_END31_I_I_I:.*]] ], [ [[TMP2]], %[[IF_THEN5_I_I]] ]
+// APPROX-NEXT: [[__R_0_I29_I_I11:%.*]] = phi i64 [ [[ADD28_I_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ 0, %[[IF_THEN5_I_I]] ]
+// APPROX-NEXT: [[__TAGP_ADDR_0_I28_I_I10:%.*]] = phi ptr [ [[INCDEC_PTR_I34_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ [[INCDEC_PTR_I_I]], %[[IF_THEN5_I_I]] ]
+// APPROX-NEXT: [[TMP4:%.*]] = add i8 [[TMP3]], -48
+// APPROX-NEXT: [[OR_COND_I32_I_I:%.*]] = icmp ult i8 [[TMP4]], 10
+// APPROX-NEXT: br i1 [[OR_COND_I32_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE_I_I_I:.*]]
// APPROX: [[IF_ELSE_I_I_I]]:
-// APPROX-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
-// APPROX-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// APPROX-NEXT: [[TMP5:%.*]] = add i8 [[TMP3]], -97
+// APPROX-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
// APPROX-NEXT: br i1 [[OR_COND33_I_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE17_I_I_I:.*]]
// APPROX: [[IF_ELSE17_I_I_I]]:
-// APPROX-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
-// APPROX-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// APPROX-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[CLEANUP_I36_I_I]]
+// APPROX-NEXT: [[TMP6:%.*]] = add i8 [[TMP3]], -65
+// APPROX-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
+// APPROX-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[_ZL3NANPKC_EXIT]]
// APPROX: [[IF_END31_I_I_I]]:
-// APPROX-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I34_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
-// APPROX-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4
-// APPROX-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// APPROX-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
+// APPROX-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I29_I_I11]], 4
+// APPROX-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// APPROX-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
-// APPROX-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
-// APPROX-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I31_I_I]], i64 1
-// APPROX-NEXT: br label %[[CLEANUP_I36_I_I]]
-// APPROX: [[CLEANUP_I36_I_I]]:
-// APPROX-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], %[[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// APPROX-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// APPROX-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, %[[IF_END31_I_I_I]] ], [ false, %[[IF_ELSE17_I_I_I]] ]
-// APPROX-NEXT: br i1 [[COND_I_I_I]], label %[[WHILE_COND_I30_I_I]], label %[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP12]]
-// APPROX: [[WHILE_COND_I_I_I]]:
-// APPROX-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], %[[CLEANUP_I_I_I:.*]] ], [ [[INCDEC_PTR_I_I]], %[[IF_THEN_I_I]] ]
-// APPROX-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], %[[CLEANUP_I_I_I]] ], [ 0, %[[IF_THEN_I_I]] ]
-// APPROX-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// APPROX-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
-// APPROX-NEXT: br i1 [[CMP_NOT_I_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I_I_I:.*]]
+// APPROX-NEXT: [[ADD28_I_I_I]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I34_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I28_I_I10]], i64 1
+// APPROX-NEXT: [[TMP7]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// APPROX-NEXT: [[CMP_NOT_I30_I_I:%.*]] = icmp eq i8 [[TMP7]], 0
+// APPROX-NEXT: br i1 [[CMP_NOT_I30_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I31_I_I]], !llvm.loop [[LOOP12]]
// APPROX: [[WHILE_BODY_I_I_I]]:
-// APPROX-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// APPROX-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
-// APPROX-NEXT: br i1 [[OR_COND_I_I_I]], label %[[IF_THEN_I_I_I:.*]], label %[[CLEANUP_I_I_I]]
+// APPROX-NEXT: [[TMP8:%.*]] = phi i8 [ [[TMP10:%.*]], %[[IF_THEN_I_I_I:.*]] ], [ [[TMP2]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[__R_0_I_I_I16:%.*]] = phi i64 [ [[SUB_I_I_I:%.*]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[__TAGP_ADDR_0_I_I_I15:%.*]] = phi ptr [ [[INCDEC_PTR_I_I_I:%.*]], %[[IF_THEN_I_I_I]] ], [ [[INCDEC_PTR_I_I]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[TMP9:%.*]] = and i8 [[TMP8]], -8
+// APPROX-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP9]], 48
+// APPROX-NEXT: br i1 [[OR_COND_I_I_I]], label %[[IF_THEN_I_I_I]], label %[[_ZL3NANPKC_EXIT]]
// APPROX: [[IF_THEN_I_I_I]]:
-// APPROX-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
-// APPROX-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// APPROX-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I16]], 3
+// APPROX-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
// APPROX-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
-// APPROX-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
-// APPROX-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1
-// APPROX-NEXT: br label %[[CLEANUP_I_I_I]]
-// APPROX: [[CLEANUP_I_I_I]]:
-// APPROX-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], %[[WHILE_BODY_I_I_I]] ]
-// APPROX-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_BODY_I_I_I]] ]
-// APPROX-NEXT: br i1 [[OR_COND_I_I_I]], label %[[WHILE_COND_I_I_I]], label %[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP8]]
-// APPROX: [[WHILE_COND_I14_I_I]]:
-// APPROX-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I_I:%.*]], %[[CLEANUP_I20_I_I:.*]] ], [ [[TAG]], %[[ENTRY]] ]
-// APPROX-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]], %[[CLEANUP_I20_I_I]] ], [ 0, %[[ENTRY]] ]
-// APPROX-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// APPROX-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// APPROX-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I18_I_I:.*]]
+// APPROX-NEXT: [[SUB_I_I_I]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I_I_I15]], i64 1
+// APPROX-NEXT: [[TMP10]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// APPROX-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP10]], 0
+// APPROX-NEXT: br i1 [[CMP_NOT_I_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I_I_I]], !llvm.loop [[LOOP8]]
// APPROX: [[WHILE_BODY_I18_I_I]]:
-// APPROX-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// APPROX-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// APPROX-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[IF_THEN_I24_I_I:.*]], label %[[CLEANUP_I20_I_I]]
-// APPROX: [[IF_THEN_I24_I_I]]:
-// APPROX-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
-// APPROX-NEXT: [[CONV5_I26_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
-// APPROX-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48
-// APPROX-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]], [[CONV5_I26_I_I]]
-// APPROX-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I_I]], i64 1
-// APPROX-NEXT: br label %[[CLEANUP_I20_I_I]]
-// APPROX: [[CLEANUP_I20_I_I]]:
-// APPROX-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [ [[INCDEC_PTR_I29_I_I]], %[[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]], %[[WHILE_BODY_I18_I_I]] ]
-// APPROX-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]], %[[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_BODY_I18_I_I]] ]
-// APPROX-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[WHILE_COND_I14_I_I]], label %[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP11]]
+// APPROX-NEXT: [[TMP11:%.*]] = phi i8 [ [[TMP13:%.*]], %[[IF_THEN_I21_I_I:.*]] ], [ [[TMP1]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[__R_0_I16_I_I7:%.*]] = phi i64 [ [[SUB_I25_I_I:%.*]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[__TAGP_ADDR_0_I15_I_I6:%.*]] = phi ptr [ [[INCDEC_PTR_I26_I_I:%.*]], %[[IF_THEN_I21_I_I]] ], [ [[TAG]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// APPROX-NEXT: [[TMP12:%.*]] = add i8 [[TMP11]], -48
+// APPROX-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP12]], 10
+// APPROX-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[IF_THEN_I21_I_I]], label %[[_ZL3NANPKC_EXIT]]
+// APPROX: [[IF_THEN_I21_I_I]]:
+// APPROX-NEXT: [[MUL_I22_I_I:%.*]] = mul i64 [[__R_0_I16_I_I7]], 10
+// APPROX-NEXT: [[CONV5_I23_I_I:%.*]] = zext nneg i8 [[TMP11]] to i64
+// APPROX-NEXT: [[ADD_I24_I_I:%.*]] = add i64 [[MUL_I22_I_I]], -48
+// APPROX-NEXT: [[SUB_I25_I_I]] = add i64 [[ADD_I24_I_I]], [[CONV5_I23_I_I]]
+// APPROX-NEXT: [[INCDEC_PTR_I26_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I_I6]], i64 1
+// APPROX-NEXT: [[TMP13]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// APPROX-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP13]], 0
+// APPROX-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I18_I_I]], !llvm.loop [[LOOP11]]
// APPROX: [[_ZL3NANPKC_EXIT]]:
-// APPROX-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ [[__R_0_I32_I_I]], %[[WHILE_COND_I30_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_COND_I_I_I]] ], [ 0, %[[CLEANUP_I_I_I]] ], [ 0, %[[CLEANUP_I36_I_I]] ], [ 0, %[[CLEANUP_I20_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_COND_I14_I_I]] ]
+// APPROX-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
// APPROX-NEXT: [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]], 2251799813685247
// APPROX-NEXT: [[BF_SET9_I:%.*]] = or disjoint i64 [[BF_VALUE_I]], 9221120237041090560
-// APPROX-NEXT: [[TMP10:%.*]] = bitcast i64 [[BF_SET9_I]] to double
-// APPROX-NEXT: ret double [[TMP10]]
+// APPROX-NEXT: [[TMP14:%.*]] = bitcast i64 [[BF_SET9_I]] to double
+// APPROX-NEXT: ret double [[TMP14]]
//
// NCRDIV-LABEL: define dso_local double @test_nan(
// NCRDIV-SAME: ptr noundef readonly captures(none) [[TAG:%.*]]) local_unnamed_addr #[[ATTR2]] {
-// NCRDIV-NEXT: [[ENTRY:.*]]:
+// NCRDIV-NEXT: [[ENTRY:.*:]]
// NCRDIV-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG]], align 1, !tbaa [[CHAR_TBAA7]]
// NCRDIV-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// NCRDIV-NEXT: br i1 [[CMP_I_I]], label %[[IF_THEN_I_I:.*]], label %[[WHILE_COND_I14_I_I:.*]]
+// NCRDIV-NEXT: br i1 [[CMP_I_I]], label %[[IF_THEN_I_I:.*]], label %[[WHILE_COND_I14_I_I_PREHEADER:.*]]
+// NCRDIV: [[WHILE_COND_I14_I_I_PREHEADER]]:
+// NCRDIV-NEXT: [[TMP1:%.*]] = load i8, ptr [[TAG]], align 1, !tbaa [[CHAR_TBAA7]]
+// NCRDIV-NEXT: [[CMP_NOT_I17_I_I5:%.*]] = icmp eq i8 [[TMP1]], 0
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I17_I_I5]], label %[[_ZL3NANPKC_EXIT:.*]], label %[[WHILE_BODY_I18_I_I:.*]]
// NCRDIV: [[IF_THEN_I_I]]:
// NCRDIV-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[TAG]], i64 1
-// NCRDIV-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// NCRDIV-NEXT: switch i8 [[TMP1]], label %[[WHILE_COND_I_I_I:.*]] [
-// NCRDIV-NEXT: i8 120, label %[[WHILE_COND_I30_I_I_PREHEADER:.*]]
-// NCRDIV-NEXT: i8 88, label %[[WHILE_COND_I30_I_I_PREHEADER]]
+// NCRDIV-NEXT: [[TMP2:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// NCRDIV-NEXT: [[CMP_NOT_I_I_I14:%.*]] = icmp eq i8 [[TMP2]], 0
+// NCRDIV-NEXT: switch i8 [[TMP2]], label %[[WHILE_COND_I_I_I_PREHEADER:.*]] [
+// NCRDIV-NEXT: i8 120, label %[[IF_THEN5_I_I:.*]]
+// NCRDIV-NEXT: i8 88, label %[[IF_THEN5_I_I]]
// NCRDIV-NEXT: ]
-// NCRDIV: [[WHILE_COND_I30_I_I_PREHEADER]]:
-// NCRDIV-NEXT: br label %[[WHILE_COND_I30_I_I:.*]]
-// NCRDIV: [[WHILE_COND_I30_I_I]]:
-// NCRDIV-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I_I:%.*]], %[[CLEANUP_I36_I_I:.*]] ], [ [[INCDEC_PTR_I_I]], %[[WHILE_COND_I30_I_I_PREHEADER]] ]
-// NCRDIV-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], %[[CLEANUP_I36_I_I]] ], [ 0, %[[WHILE_COND_I30_I_I_PREHEADER]] ]
-// NCRDIV-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// NCRDIV-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// NCRDIV-NEXT: br i1 [[CMP_NOT_I33_I_I]], label %[[_ZL3NANPKC_EXIT:.*]], label %[[WHILE_BODY_I34_I_I:.*]]
-// NCRDIV: [[WHILE_BODY_I34_I_I]]:
-// NCRDIV-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// NCRDIV-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// NCRDIV-NEXT: br i1 [[OR_COND_I35_I_I]], label %[[IF_END31_I_I_I:.*]], label %[[IF_ELSE_I_I_I:.*]]
+// NCRDIV: [[WHILE_COND_I_I_I_PREHEADER]]:
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I_I_I14]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I_I_I:.*]]
+// NCRDIV: [[IF_THEN5_I_I]]:
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I_I_I14]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I31_I_I:.*]]
+// NCRDIV: [[WHILE_BODY_I31_I_I]]:
+// NCRDIV-NEXT: [[TMP3:%.*]] = phi i8 [ [[TMP7:%.*]], %[[IF_END31_I_I_I:.*]] ], [ [[TMP2]], %[[IF_THEN5_I_I]] ]
+// NCRDIV-NEXT: [[__R_0_I29_I_I11:%.*]] = phi i64 [ [[ADD28_I_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ 0, %[[IF_THEN5_I_I]] ]
+// NCRDIV-NEXT: [[__TAGP_ADDR_0_I28_I_I10:%.*]] = phi ptr [ [[INCDEC_PTR_I34_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ [[INCDEC_PTR_I_I]], %[[IF_THEN5_I_I]] ]
+// NCRDIV-NEXT: [[TMP4:%.*]] = add i8 [[TMP3]], -48
+// NCRDIV-NEXT: [[OR_COND_I32_I_I:%.*]] = icmp ult i8 [[TMP4]], 10
+// NCRDIV-NEXT: br i1 [[OR_COND_I32_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE_I_I_I:.*]]
// NCRDIV: [[IF_ELSE_I_I_I]]:
-// NCRDIV-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
-// NCRDIV-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// NCRDIV-NEXT: [[TMP5:%.*]] = add i8 [[TMP3]], -97
+// NCRDIV-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
// NCRDIV-NEXT: br i1 [[OR_COND33_I_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE17_I_I_I:.*]]
// NCRDIV: [[IF_ELSE17_I_I_I]]:
-// NCRDIV-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
-// NCRDIV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// NCRDIV-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[CLEANUP_I36_I_I]]
+// NCRDIV-NEXT: [[TMP6:%.*]] = add i8 [[TMP3]], -65
+// NCRDIV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
+// NCRDIV-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[_ZL3NANPKC_EXIT]]
// NCRDIV: [[IF_END31_I_I_I]]:
-// NCRDIV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I34_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
-// NCRDIV-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4
-// NCRDIV-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// NCRDIV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
+// NCRDIV-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I29_I_I11]], 4
+// NCRDIV-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// NCRDIV-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
-// NCRDIV-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
-// NCRDIV-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I31_I_I]], i64 1
-// NCRDIV-NEXT: br label %[[CLEANUP_I36_I_I]]
-// NCRDIV: [[CLEANUP_I36_I_I]]:
-// NCRDIV-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], %[[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// NCRDIV-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// NCRDIV-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, %[[IF_END31_I_I_I]] ], [ false, %[[IF_ELSE17_I_I_I]] ]
-// NCRDIV-NEXT: br i1 [[COND_I_I_I]], label %[[WHILE_COND_I30_I_I]], label %[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP12]]
-// NCRDIV: [[WHILE_COND_I_I_I]]:
-// NCRDIV-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]] ]
-// NCRDIV-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], %[[CLEANUP_I_I_I]] ], [ 0, %[[IF_THEN_I_I]] ]
-// NCRDIV-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// NCRDIV-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
-// NCRDIV-NEXT: br i1 [[CMP_NOT_I_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I_I_I:.*]]
+// NCRDIV-NEXT: [[ADD28_I_I_I]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I34_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I28_I_I10]], i64 1
+// NCRDIV-NEXT: [[TMP7]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// NCRDIV-NEXT: [[CMP_NOT_I30_I_I:%.*]] = icmp eq i8 [[TMP7]], 0
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I30_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I31_I_I]], !llvm.loop [[LOOP12]]
// NCRDIV: [[WHILE_BODY_I_I_I]]:
-// NCRDIV-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// NCRDIV-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
-// NCRDIV-NEXT: br i1 [[OR_COND_I_I_I]], label %[[IF_THEN_I_I_I:.*]], label %[[CLEANUP_I_I_I]]
+// NCRDIV-NEXT: [[TMP8:%.*]] = phi i8 [ [[TMP10:%.*]], %[[IF_THEN_I_I_I:.*]] ], [ [[TMP2]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[__R_0_I_I_I16:%.*]] = phi i64 [ [[SUB_I_I_I:%.*]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[__TAGP_ADDR_0_I_I_I15:%.*]] = phi ptr [ [[INCDEC_PTR_I_I_I:%.*]], %[[IF_THEN_I_I_I]] ], [ [[INCDEC_PTR_I_I]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[TMP9:%.*]] = and i8 [[TMP8]], -8
+// NCRDIV-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP9]], 48
+// NCRDIV-NEXT: br i1 [[OR_COND_I_I_I]], label %[[IF_THEN_I_I_I]], label %[[_ZL3NANPKC_EXIT]]
// NCRDIV: [[IF_THEN_I_I_I]]:
-// NCRDIV-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
-// NCRDIV-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// NCRDIV-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I16]], 3
+// NCRDIV-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
// NCRDIV-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
-// NCRDIV-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
-// NCRDIV-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1
-// NCRDIV-NEXT: br label %[[CLEANUP_I_I_I]]
-// NCRDIV: [[CLEANUP_I_I_I]]:
-// NCRDIV-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]] ]
-// NCRDIV-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]] ]
-// NCRDIV-NEXT: br i1 [[OR_COND_I_I_I]], label %[[WHILE_COND_I_I_I]], label %[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP8]]
-// NCRDIV: [[WHILE_COND_I14_I_I]]:
-// NCRDIV-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I_I:%.*]], %[[CLEANUP_I20_I_I:.*]] ], [ [[TAG]], %[[ENTRY]] ]
-// NCRDIV-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]], %[[CLEANUP_I20_I_I]] ], [ 0, %[[ENTRY]] ]
-// NCRDIV-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
-// NCRDIV-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// NCRDIV-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I18_I_I:.*]]
+// NCRDIV-NEXT: [[SUB_I_I_I]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I_I_I15]], i64 1
+// NCRDIV-NEXT: [[TMP10]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// NCRDIV-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP10]], 0
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I_I_I]], !llvm.loop [[LOOP8]]
// NCRDIV: [[WHILE_BODY_I18_I_I]]:
-// NCRDIV-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// NCRDIV-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// NCRDIV-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[IF_THEN_I24_I_I:.*]], label %[[CLEANUP_I20_I_I]]
-// NCRDIV: [[IF_THEN_I24_I_I]]:
-// NCRDIV-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
-// NCRDIV-NEXT: [[CONV5_I26_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
-// NCRDIV-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48
-// NCRDIV-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]], [[CONV5_I26_I_I]]
-// NCRDIV-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I_I]], i64 1
-// NCRDIV-NEXT: br label %[[CLEANUP_I20_I_I]]
-// NCRDIV: [[CLEANUP_I20_I_I]]:
-// NCRDIV-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [ [[INCDEC_PTR_I29_I_I]], %[[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]], %[[WHILE_BODY_I18_I_I]] ]
-// NCRDIV-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]], %[[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_BODY_I18_I_I]] ]
-// NCRDIV-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[WHILE_COND_I14_I_I]], label %[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP11]]
+// NCRDIV-NEXT: [[TMP11:%.*]] = phi i8 [ [[TMP13:%.*]], %[[IF_THEN_I21_I_I:.*]] ], [ [[TMP1]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[__R_0_I16_I_I7:%.*]] = phi i64 [ [[SUB_I25_I_I:%.*]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[__TAGP_ADDR_0_I15_I_I6:%.*]] = phi ptr [ [[INCDEC_PTR_I26_I_I:%.*]], %[[IF_THEN_I21_I_I]] ], [ [[TAG]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// NCRDIV-NEXT: [[TMP12:%.*]] = add i8 [[TMP11]], -48
+// NCRDIV-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP12]], 10
+// NCRDIV-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[IF_THEN_I21_I_I]], label %[[_ZL3NANPKC_EXIT]]
+// NCRDIV: [[IF_THEN_I21_I_I]]:
+// NCRDIV-NEXT: [[MUL_I22_I_I:%.*]] = mul i64 [[__R_0_I16_I_I7]], 10
+// NCRDIV-NEXT: [[CONV5_I23_I_I:%.*]] = zext nneg i8 [[TMP11]] to i64
+// NCRDIV-NEXT: [[ADD_I24_I_I:%.*]] = add i64 [[MUL_I22_I_I]], -48
+// NCRDIV-NEXT: [[SUB_I25_I_I]] = add i64 [[ADD_I24_I_I]], [[CONV5_I23_I_I]]
+// NCRDIV-NEXT: [[INCDEC_PTR_I26_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I15_I_I6]], i64 1
+// NCRDIV-NEXT: [[TMP13]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA7]]
+// NCRDIV-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP13]], 0
+// NCRDIV-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I18_I_I]], !llvm.loop [[LOOP11]]
// NCRDIV: [[_ZL3NANPKC_EXIT]]:
-// NCRDIV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ [[__R_0_I32_I_I]], %[[WHILE_COND_I30_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_COND_I_I_I]] ], [ 0, %[[CLEANUP_I_I_I]] ], [ 0, %[[CLEANUP_I36_I_I]] ], [ 0, %[[CLEANUP_I20_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_COND_I14_I_I]] ]
+// NCRDIV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
// NCRDIV-NEXT: [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]], 2251799813685247
// NCRDIV-NEXT: [[BF_SET9_I:%.*]] = or disjoint i64 [[BF_VALUE_I]], 9221120237041090560
-// NCRDIV-NEXT: [[TMP10:%.*]] = bitcast i64 [[BF_SET9_I]] to double
-// NCRDIV-NEXT: ret double [[TMP10]]
+// NCRDIV-NEXT: [[TMP14:%.*]] = bitcast i64 [[BF_SET9_I]] to double
+// NCRDIV-NEXT: ret double [[TMP14]]
//
// AMDGCNSPIRV-LABEL: define spir_func double @test_nan(
// AMDGCNSPIRV-SAME: ptr addrspace(4) noundef readonly captures(none) [[TAG:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR2]] {
@@ -5032,53 +4947,49 @@ extern "C" __device__ float test_nanf(const char *tag) {
// AMDGCNSPIRV-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[TAG]], i64 1
// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = load i8, ptr addrspace(4) [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
// AMDGCNSPIRV-NEXT: switch i8 [[TMP1]], label %[[WHILE_COND_I_I_I:.*]] [
-// AMDGCNSPIRV-NEXT: i8 120, label %[[WHILE_COND_I28_I_I_PREHEADER:.*]]
-// AMDGCNSPIRV-NEXT: i8 88, label %[[WHILE_COND_I28_I_I_PREHEADER]]
+// AMDGCNSPIRV-NEXT: i8 120, label %[[IF_THEN5_I_I:.*]]
+// AMDGCNSPIRV-NEXT: i8 88, label %[[IF_THEN5_I_I]]
// AMDGCNSPIRV-NEXT: ]
-// AMDGCNSPIRV: [[WHILE_COND_I28_I_I_PREHEADER]]:
-// AMDGCNSPIRV-NEXT: br label %[[WHILE_COND_I28_I_I:.*]]
-// AMDGCNSPIRV: [[WHILE_COND_I28_I_I]]:
-// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_0_I29_I_I:%.*]] = phi ptr addrspace(4) [ [[__TAGP_ADDR_1_I34_I_I:%.*]], %[[CLEANUP_I_I_I:.*]] ], [ [[INCDEC_PTR_I_I]], %[[WHILE_COND_I28_I_I_PREHEADER]] ]
-// AMDGCNSPIRV-NEXT: [[__R_0_I30_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], %[[CLEANUP_I_I_I]] ], [ 0, %[[WHILE_COND_I28_I_I_PREHEADER]] ]
-// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I29_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// AMDGCNSPIRV-NEXT: [[CMP_NOT_I31_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I31_I_I]], label %[[_ZL3NANPKC_EXIT:.*]], label %[[WHILE_BODY_I32_I_I:.*]]
+// AMDGCNSPIRV: [[IF_THEN5_I_I]]:
+// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = load i8, ptr addrspace(4) [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// AMDGCNSPIRV-NEXT: [[CMP_NOT_I31_I_I5:%.*]] = icmp eq i8 [[TMP2]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I31_I_I5]], label %[[_ZL3NANPKC_EXIT:.*]], label %[[WHILE_BODY_I32_I_I:.*]]
// AMDGCNSPIRV: [[WHILE_BODY_I32_I_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// AMDGCNSPIRV-NEXT: [[OR_COND_I33_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// AMDGCNSPIRV-NEXT: br i1 [[OR_COND_I33_I_I]], label %[[IF_END31_I_I_I:.*]], label %[[IF_ELSE_I_I_I:.*]]
+// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = phi i8 [ [[TMP7:%.*]], %[[IF_END31_I_I_I:.*]] ], [ [[TMP2]], %[[IF_THEN5_I_I]] ]
+// AMDGCNSPIRV-NEXT: [[__R_0_I30_I_I7:%.*]] = phi i64 [ [[ADD28_I_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ 0, %[[IF_THEN5_I_I]] ]
+// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_0_I29_I_I6:%.*]] = phi ptr addrspace(4) [ [[INCDEC_PTR_I36_I_I:%.*]], %[[IF_END31_I_I_I]] ], [ [[INCDEC_PTR_I_I]], %[[IF_THEN5_I_I]] ]
+// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = add i8 [[TMP3]], -48
+// AMDGCNSPIRV-NEXT: [[OR_COND_I33_I_I:%.*]] = icmp ult i8 [[TMP4]], 10
+// AMDGCNSPIRV-NEXT: br i1 [[OR_COND_I33_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE_I_I_I:.*]]
// AMDGCNSPIRV: [[IF_ELSE_I_I_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
-// AMDGCNSPIRV-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
+// AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = add i8 [[TMP3]], -97
+// AMDGCNSPIRV-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
// AMDGCNSPIRV-NEXT: br i1 [[OR_COND33_I_I_I]], label %[[IF_END31_I_I_I]], label %[[IF_ELSE17_I_I_I:.*]]
// AMDGCNSPIRV: [[IF_ELSE17_I_I_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
-// AMDGCNSPIRV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// AMDGCNSPIRV-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[CLEANUP_I_I_I]]
+// AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = add i8 [[TMP3]], -65
+// AMDGCNSPIRV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
+// AMDGCNSPIRV-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[_ZL3NANPKC_EXIT]]
// AMDGCNSPIRV: [[IF_END31_I_I_I]]:
// AMDGCNSPIRV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I32_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
-// AMDGCNSPIRV-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I30_I_I]], 4
-// AMDGCNSPIRV-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP2]] to i64
+// AMDGCNSPIRV-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I30_I_I7]], 4
+// AMDGCNSPIRV-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// AMDGCNSPIRV-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
-// AMDGCNSPIRV-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
-// AMDGCNSPIRV-NEXT: [[INCDEC_PTR_I37_I_I:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__TAGP_ADDR_0_I29_I_I]], i64 1
-// AMDGCNSPIRV-NEXT: br label %[[CLEANUP_I_I_I]]
-// AMDGCNSPIRV: [[CLEANUP_I_I_I]]:
-// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_1_I34_I_I]] = phi ptr addrspace(4) [ [[INCDEC_PTR_I37_I_I]], %[[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I29_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// AMDGCNSPIRV-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ [[__R_0_I30_I_I]], %[[IF_ELSE17_I_I_I]] ]
-// AMDGCNSPIRV-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, %[[IF_END31_I_I_I]] ], [ false, %[[IF_ELSE17_I_I_I]] ]
-// AMDGCNSPIRV-NEXT: br i1 [[COND_I_I_I]], label %[[WHILE_COND_I28_I_I]], label %[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP13]]
+// AMDGCNSPIRV-NEXT: [[ADD28_I_I_I]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
+// AMDGCNSPIRV-NEXT: [[INCDEC_PTR_I36_I_I]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__TAGP_ADDR_0_I29_I_I6]], i64 1
+// AMDGCNSPIRV-NEXT: [[TMP7]] = load i8, ptr addrspace(4) [[INCDEC_PTR_I36_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// AMDGCNSPIRV-NEXT: [[CMP_NOT_I31_I_I:%.*]] = icmp eq i8 [[TMP7]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I31_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I32_I_I]], !llvm.loop [[LOOP13]]
// AMDGCNSPIRV: [[WHILE_COND_I_I_I]]:
// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr addrspace(4) [ [[__TAGP_ADDR_1_I_I_I:%.*]], %[[WHILE_BODY_I_I_I:.*]] ], [ [[INCDEC_PTR_I_I]], %[[IF_THEN_I_I]] ]
// AMDGCNSPIRV-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_THEN_I_I]] ]
-// AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// AMDGCNSPIRV-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// AMDGCNSPIRV-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I_I_I]]
// AMDGCNSPIRV: [[WHILE_BODY_I_I_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// AMDGCNSPIRV-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = and i8 [[TMP8]], -8
+// AMDGCNSPIRV-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP9]], 48
// AMDGCNSPIRV-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
-// AMDGCNSPIRV-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP6]] to i64
+// AMDGCNSPIRV-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
// AMDGCNSPIRV-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
// AMDGCNSPIRV-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_1_I_I_I_IDX:%.*]] = zext i1 [[OR_COND_I_I_I]] to i64
@@ -5088,14 +4999,14 @@ extern "C" __device__ float test_nanf(const char *tag) {
// AMDGCNSPIRV: [[WHILE_COND_I14_I_I]]:
// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr addrspace(4) [ [[__TAGP_ADDR_1_I25_I_I:%.*]], %[[WHILE_BODY_I18_I_I:.*]] ], [ [[TAG]], %[[ENTRY]] ]
// AMDGCNSPIRV-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I26_I_I:%.*]], %[[WHILE_BODY_I18_I_I]] ], [ 0, %[[ENTRY]] ]
-// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// AMDGCNSPIRV-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
+// AMDGCNSPIRV-NEXT: [[TMP10:%.*]] = load i8, ptr addrspace(4) [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// AMDGCNSPIRV-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP10]], 0
// AMDGCNSPIRV-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I18_I_I]]
// AMDGCNSPIRV: [[WHILE_BODY_I18_I_I]]:
-// AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// AMDGCNSPIRV-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
+// AMDGCNSPIRV-NEXT: [[TMP11:%.*]] = add i8 [[TMP10]], -48
+// AMDGCNSPIRV-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP11]], 10
// AMDGCNSPIRV-NEXT: [[MUL_I20_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10
-// AMDGCNSPIRV-NEXT: [[CONV5_I21_I_I:%.*]] = zext nneg i8 [[TMP8]] to i64
+// AMDGCNSPIRV-NEXT: [[CONV5_I21_I_I:%.*]] = zext nneg i8 [[TMP10]] to i64
// AMDGCNSPIRV-NEXT: [[ADD_I22_I_I:%.*]] = add i64 [[MUL_I20_I_I]], -48
// AMDGCNSPIRV-NEXT: [[SUB_I23_I_I:%.*]] = add i64 [[ADD_I22_I_I]], [[CONV5_I21_I_I]]
// AMDGCNSPIRV-NEXT: [[__TAGP_ADDR_1_I25_I_I_IDX:%.*]] = zext i1 [[OR_COND_I19_I_I]] to i64
@@ -5103,11 +5014,11 @@ extern "C" __device__ float test_nanf(const char *tag) {
// AMDGCNSPIRV-NEXT: [[__R_1_I26_I_I]] = select i1 [[OR_COND_I19_I_I]], i64 [[SUB_I23_I_I]], i64 [[__R_0_I16_I_I]]
// AMDGCNSPIRV-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[WHILE_COND_I14_I_I]], label %[[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP12]]
// AMDGCNSPIRV: [[_ZL3NANPKC_EXIT]]:
-// AMDGCNSPIRV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ [[__R_0_I30_I_I]], %[[WHILE_COND_I28_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_COND_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[CLEANUP_I_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_COND_I14_I_I]] ]
+// AMDGCNSPIRV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_COND_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_COND_I14_I_I]] ]
// AMDGCNSPIRV-NEXT: [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]], 2251799813685247
// AMDGCNSPIRV-NEXT: [[BF_SET9_I:%.*]] = or disjoint i64 [[BF_VALUE_I]], 9221120237041090560
-// AMDGCNSPIRV-NEXT: [[TMP10:%.*]] = bitcast i64 [[BF_SET9_I]] to double
-// AMDGCNSPIRV-NEXT: ret double [[TMP10]]
+// AMDGCNSPIRV-NEXT: [[TMP12:%.*]] = bitcast i64 [[BF_SET9_I]] to double
+// AMDGCNSPIRV-NEXT: ret double [[TMP12]]
//
extern "C" __device__ double test_nan(const char *tag) {
return nan(tag);
diff --git a/llvm/lib/Transforms/Utils/SimplifyCFG.cpp b/llvm/lib/Transforms/Utils/SimplifyCFG.cpp
index f1d47abe79365..a8fd9f4a10064 100644
--- a/llvm/lib/Transforms/Utils/SimplifyCFG.cpp
+++ b/llvm/lib/Transforms/Utils/SimplifyCFG.cpp
@@ -201,6 +201,11 @@ static cl::opt<unsigned> MaxSwitchCasesPerResult(
"max-switch-cases-per-result", cl::Hidden, cl::init(16),
cl::desc("Limit cases to analyze when converting a switch to select"));
+static cl::opt<unsigned> MaxJumpThreadingLiveBlocks(
+ "max-jump-threading-live-blocks", cl::Hidden, cl::init(24),
+ cl::desc("Limit number of blocks a define in a threaded block is allowed "
+ "to be live in"));
+
extern cl::opt<bool> ProfcheckDisableMetadataFixes;
} // end namespace llvm
@@ -3439,8 +3444,27 @@ bool SimplifyCFGOpt::speculativelyExecuteBB(CondBrInst *BI,
return true;
}
+using BlocksSet = SmallPtrSet<BasicBlock *, 8>;
+
+// Return false if number of blocks searched is too much.
+static bool findReaching(BasicBlock *BB, BasicBlock *DefBB,
+ BlocksSet &ReachesNonLocalUses) {
+ if (BB == DefBB)
+ return true;
+ if (!ReachesNonLocalUses.insert(BB).second)
+ return true;
+
+ if (ReachesNonLocalUses.size() > MaxJumpThreadingLiveBlocks)
+ return false;
+ for (BasicBlock *Pred : predecessors(BB))
+ if (!findReaching(Pred, DefBB, ReachesNonLocalUses))
+ return false;
+ return true;
+}
+
/// Return true if we can thread a branch across this block.
-static bool blockIsSimpleEnoughToThreadThrough(BasicBlock *BB) {
+static bool blockIsSimpleEnoughToThreadThrough(BasicBlock *BB,
+ BlocksSet &NonLocalUseBlocks) {
int Size = 0;
EphemeralValueTracker EphTracker;
@@ -3460,12 +3484,16 @@ static bool blockIsSimpleEnoughToThreadThrough(BasicBlock *BB) {
return false; // Don't clone large BB's.
}
- // We can only support instructions that do not define values that are
- // live outside of the current basic block.
+ // Record blocks with non-local uses of values defined in the current basic
+ // block.
for (User *U : I.users()) {
Instruction *UI = cast<Instruction>(U);
- if (UI->getParent() != BB || isa<PHINode>(UI))
- return false;
+ BasicBlock *UsedInBB = UI->getParent();
+ if (UsedInBB == BB) {
+ if (isa<PHINode>(UI))
+ return false;
+ } else
+ NonLocalUseBlocks.insert(UsedInBB);
}
// Looks ok, continue checking.
@@ -3524,18 +3552,37 @@ foldCondBranchOnValueKnownInPredecessorImpl(CondBrInst *BI, DomTreeUpdater *DTU,
return false;
// Now we know that this block has multiple preds and two succs.
- // Check that the block is small enough and values defined in the block are
- // not used outside of it.
- if (!blockIsSimpleEnoughToThreadThrough(BB))
+ // Check that the block is small enough and record which non-local blocks use
+ // values defined in the block.
+
+ BlocksSet NonLocalUseBlocks;
+ BlocksSet ReachesNonLocalUseBlocks;
+ if (!blockIsSimpleEnoughToThreadThrough(BB, NonLocalUseBlocks))
return false;
+ // Jump-threading can only be done to destinations where no values defined
+ // in BB are live.
+
+ // Quickly check if both destinations have uses. If so, jump-threading cannot
+ // be done.
+ if (NonLocalUseBlocks.contains(BI->getSuccessor(0)) &&
+ NonLocalUseBlocks.contains(BI->getSuccessor(1)))
+ return false;
+
+ // Search backward from NonLocalUseBlocks to find which blocks
+ // reach non-local uses.
+ for (BasicBlock *UseBB : NonLocalUseBlocks)
+ // Give up if too many blocks are searched.
+ if (!findReaching(UseBB, BB, ReachesNonLocalUseBlocks))
+ return false;
+
for (const auto &Pair : KnownValues) {
- // Okay, we now know that all edges from PredBB should be revectored to
- // branch to RealDest.
ConstantInt *CB = Pair.first;
ArrayRef<BasicBlock *> PredBBs = Pair.second.getArrayRef();
BasicBlock *RealDest = BI->getSuccessor(!CB->getZExtValue());
+ // Okay, we now know that all edges from PredBB should be revectored to
+ // branch to RealDest.
if (RealDest == BB)
continue; // Skip self loops.
@@ -3545,6 +3592,10 @@ foldCondBranchOnValueKnownInPredecessorImpl(CondBrInst *BI, DomTreeUpdater *DTU,
}))
continue;
+ // Only revector to RealDest if no values defined in BB are live.
+ if (ReachesNonLocalUseBlocks.contains(RealDest))
+ continue;
+
LLVM_DEBUG({
dbgs() << "Condition " << *Cond << " in " << BB->getName()
<< " has value " << *Pair.first << " in predecessors:\n";
diff --git a/llvm/test/CodeGen/AArch64/avoid-free-ext-promotion.ll b/llvm/test/CodeGen/AArch64/avoid-free-ext-promotion.ll
index 634d1b90ff903..5f5b27af5c8cf 100644
--- a/llvm/test/CodeGen/AArch64/avoid-free-ext-promotion.ll
+++ b/llvm/test/CodeGen/AArch64/avoid-free-ext-promotion.ll
@@ -59,37 +59,33 @@ bb27: ; preds = %bb9, %bb8
define void @avoid_promotion_2_and(ptr nocapture noundef %arg) {
; CHECK-LABEL: avoid_promotion_2_and:
; CHECK: ; %bb.0: ; %entry
-; CHECK-NEXT: add x8, x0, #32
-; CHECK-NEXT: b LBB1_2
-; CHECK-NEXT: LBB1_1: ; %latch
-; CHECK-NEXT: ; in Loop: Header=BB1_2 Depth=1
-; CHECK-NEXT: cmp w9, #2
-; CHECK-NEXT: add x8, x8, #56
-; CHECK-NEXT: b.ls LBB1_4
-; CHECK-NEXT: LBB1_2: ; %loop
+; CHECK-NEXT: mov x8, xzr
+; CHECK-NEXT: add x9, x0, #32
+; CHECK-NEXT: LBB1_1: ; %loop
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
-; CHECK-NEXT: ldr w9, [x8, #20]
-; CHECK-NEXT: cmp w9, #3
-; CHECK-NEXT: b.lo LBB1_1
-; CHECK-NEXT: ; %bb.3: ; %then
-; CHECK-NEXT: ; in Loop: Header=BB1_2 Depth=1
-; CHECK-NEXT: ldp w13, w12, [x8, #12]
-; CHECK-NEXT: ldr w10, [x8]
+; CHECK-NEXT: ldr w10, [x9, #20]
+; CHECK-NEXT: cmp w10, #3
+; CHECK-NEXT: b.lo LBB1_3
+; CHECK-NEXT: ; %bb.2: ; %then
+; CHECK-NEXT: ; in Loop: Header=BB1_1 Depth=1
+; CHECK-NEXT: ldp w13, w12, [x9, #12]
+; CHECK-NEXT: ldr w10, [x9]
; CHECK-NEXT: ldr x11, [x0]
-; CHECK-NEXT: ldr w14, [x8, #8]
+; CHECK-NEXT: add x8, x8, #1
+; CHECK-NEXT: ldr w14, [x9, #8]
; CHECK-NEXT: lsl w10, w10, w13
; CHECK-NEXT: ldrb w11, [x11, x12]
; CHECK-NEXT: eor w10, w10, w11
-; CHECK-NEXT: ldur w11, [x8, #-24]
+; CHECK-NEXT: ldur w11, [x9, #-24]
; CHECK-NEXT: and w10, w10, w14
-; CHECK-NEXT: ldp x14, x13, [x8, #-16]
-; CHECK-NEXT: str w10, [x8]
+; CHECK-NEXT: ldp x14, x13, [x9, #-16]
+; CHECK-NEXT: str w10, [x9], #56
; CHECK-NEXT: and w11, w11, w12
; CHECK-NEXT: ldrh w15, [x13, w10, uxtw #1]
; CHECK-NEXT: strh w15, [x14, w11, uxtw #1]
; CHECK-NEXT: strh w12, [x13, w10, uxtw #1]
; CHECK-NEXT: b LBB1_1
-; CHECK-NEXT: LBB1_4: ; %exit
+; CHECK-NEXT: LBB1_3: ; %exit.critedge
; CHECK-NEXT: ret
entry:
br label %loop
diff --git a/llvm/test/Transforms/SimplifyCFG/jump-threading-live-on-exit.ll b/llvm/test/Transforms/SimplifyCFG/jump-threading-live-on-exit.ll
new file mode 100644
index 0000000000000..32b77191c9304
--- /dev/null
+++ b/llvm/test/Transforms/SimplifyCFG/jump-threading-live-on-exit.ll
@@ -0,0 +1,195 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 5
+; RUN: opt -passes=simplifycfg -S < %s | FileCheck %s
+
+; Allow jump-threading when values defined in the block are live outside of the block
+; to those destinations in which the values are dead.
+
+define void @testA(ptr %ptrA, ptr %ptrB, i64 %a, i64 %b) {
+; CHECK-LABEL: define void @testA(
+; CHECK-SAME: ptr [[PTRA:%.*]], ptr [[PTRB:%.*]], i64 [[A:%.*]], i64 [[B:%.*]]) {
+; CHECK-NEXT: [[MAINA:.*:]]
+; CHECK-NEXT: [[COND:%.*]] = icmp slt i64 [[A]], [[B]]
+; CHECK-NEXT: br i1 [[COND]], label %[[IFA:.*]], label %[[MAINC:.*]]
+; CHECK: [[IFA]]:
+; CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[PTRA]], align 4
+; CHECK-NEXT: store i64 [[TMP0]], ptr [[PTRB]], align 4
+; CHECK-NEXT: br label %[[MAINC]]
+; CHECK: [[MAINC]]:
+; CHECK-NEXT: ret void
+;
+mainA:
+ %cond = icmp slt i64 %a, %b
+ br i1 %cond, label %ifA, label %mainB
+
+ifA:
+ %518 = load i64, ptr %ptrA
+ br label %mainB
+
+; %value is live outside of block mainB, but jump-threading
+; can still occur to destination mainC, since %value is dead there.
+; Subsequent CFG simplifications will create one if block.
+mainB:
+ %value = phi i64 [ %518, %ifA ], [ zeroinitializer, %mainA ]
+ br i1 %cond, label %ifB, label %mainC
+
+ifB:
+ store i64 %value, ptr %ptrB
+ br label %mainC
+
+mainC:
+ ret void
+}
+
+
+define void @testB(ptr %ptrA, ptr %ptrB, i64 %a, i64 %b, i64 %c) {
+; CHECK-LABEL: define void @testB(
+; CHECK-SAME: ptr [[PTRA:%.*]], ptr [[PTRB:%.*]], i64 [[A:%.*]], i64 [[B:%.*]], i64 [[C:%.*]]) {
+; CHECK-NEXT: [[MAINA:.*:]]
+; CHECK-NEXT: [[COND:%.*]] = icmp slt i64 [[A]], [[B]]
+; CHECK-NEXT: br i1 [[COND]], label %[[IFA:.*]], label %[[MAINC:.*]]
+; CHECK: [[IFA]]:
+; CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[PTRA]], align 4
+; CHECK-NEXT: [[COND2:%.*]] = icmp slt i64 [[A]], [[C]]
+; CHECK-NEXT: [[PTR_ARM1:%.*]] = getelementptr i64, ptr [[PTRB]], i64 8
+; CHECK-NEXT: [[PTR_ARM2:%.*]] = getelementptr i64, ptr [[PTRB]], i64 16
+; CHECK-NEXT: [[PTRC:%.*]] = select i1 [[COND2]], ptr [[PTR_ARM1]], ptr [[PTR_ARM2]]
+; CHECK-NEXT: store i64 [[TMP0]], ptr [[PTRC]], align 4
+; CHECK-NEXT: br label %[[MAINC]]
+; CHECK: [[MAINC]]:
+; CHECK-NEXT: ret void
+;
+mainA:
+ %cond = icmp slt i64 %a, %b
+ br i1 %cond, label %ifA, label %mainB
+
+ifA:
+ %518 = load i64, ptr %ptrA
+ br label %mainB
+
+; Use of %value is not in either immediate destination of mainB.
+mainB:
+ %value = phi i64 [ %518, %ifA ], [ zeroinitializer, %mainA ]
+ br i1 %cond, label %ifB, label %mainC
+
+ifB:
+ %cond2 = icmp slt i64 %a, %c
+ br i1 %cond2, label %ifB_arm1, label %ifB_arm2
+
+ifB_arm1:
+ %ptr_arm1 = getelementptr i64, ptr %ptrB, i64 8
+ br label %ifB_join
+
+ifB_arm2:
+ %ptr_arm2 = getelementptr i64, ptr %ptrB, i64 16
+ br label %ifB_join
+
+ifB_join:
+ %ptrC = phi ptr [ %ptr_arm1, %ifB_arm1 ], [ %ptr_arm2, %ifB_arm2 ]
+ store i64 %value, ptr %ptrC
+ br label %mainC
+
+mainC:
+ ret void
+}
+
+
+; Jump-threading is not done since %value is live in both destinations.
+define void @testA_negative(ptr %ptrA, ptr %ptrB, ptr %ptrD, i64 %a, i64 %b) {
+; CHECK-LABEL: define void @testA_negative(
+; CHECK-SAME: ptr [[PTRA:%.*]], ptr [[PTRB:%.*]], ptr [[PTRD:%.*]], i64 [[A:%.*]], i64 [[B:%.*]]) {
+; CHECK-NEXT: [[MAINA:.*]]:
+; CHECK-NEXT: [[COND:%.*]] = icmp slt i64 [[A]], [[B]]
+; CHECK-NEXT: br i1 [[COND]], label %[[IFA:.*]], label %[[MAINB:.*]]
+; CHECK: [[IFA]]:
+; CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[PTRA]], align 4
+; CHECK-NEXT: br label %[[MAINB]]
+; CHECK: [[MAINB]]:
+; CHECK-NEXT: [[VALUE:%.*]] = phi i64 [ [[TMP0]], %[[IFA]] ], [ 0, %[[MAINA]] ]
+; CHECK-NEXT: br i1 [[COND]], label %[[IFB:.*]], label %[[MAINC:.*]]
+; CHECK: [[IFB]]:
+; CHECK-NEXT: store i64 [[VALUE]], ptr [[PTRB]], align 4
+; CHECK-NEXT: br label %[[MAINC]]
+; CHECK: [[MAINC]]:
+; CHECK-NEXT: store i64 [[VALUE]], ptr [[PTRD]], align 4
+; CHECK-NEXT: ret void
+;
+mainA:
+ %cond = icmp slt i64 %a, %b
+ br i1 %cond, label %ifA, label %mainB
+
+ifA:
+ %518 = load i64, ptr %ptrA
+ br label %mainB
+
+mainB:
+ %value = phi i64 [ %518, %ifA ], [ zeroinitializer, %mainA ]
+ br i1 %cond, label %ifB, label %mainC
+
+ifB:
+ store i64 %value, ptr %ptrB
+ br label %mainC
+
+mainC:
+ store i64 %value, ptr %ptrD
+ ret void
+}
+
+
+; Jump-threading is not done since %value is live in both destinations.
+define void @testB_negative(ptr %ptrA, ptr %ptrB, ptr %ptrD, i64 %a, i64 %b, i64 %c) {
+; CHECK-LABEL: define void @testB_negative(
+; CHECK-SAME: ptr [[PTRA:%.*]], ptr [[PTRB:%.*]], ptr [[PTRD:%.*]], i64 [[A:%.*]], i64 [[B:%.*]], i64 [[C:%.*]]) {
+; CHECK-NEXT: [[MAINA:.*]]:
+; CHECK-NEXT: [[COND:%.*]] = icmp slt i64 [[A]], [[B]]
+; CHECK-NEXT: br i1 [[COND]], label %[[IFA:.*]], label %[[MAINB:.*]]
+; CHECK: [[IFA]]:
+; CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[PTRA]], align 4
+; CHECK-NEXT: br label %[[MAINB]]
+; CHECK: [[MAINB]]:
+; CHECK-NEXT: [[VALUE:%.*]] = phi i64 [ [[TMP0]], %[[IFA]] ], [ 0, %[[MAINA]] ]
+; CHECK-NEXT: br i1 [[COND]], label %[[IFB:.*]], label %[[MAINC:.*]]
+; CHECK: [[IFB]]:
+; CHECK-NEXT: [[COND2:%.*]] = icmp slt i64 [[A]], [[C]]
+; CHECK-NEXT: [[PTR_ARM1:%.*]] = getelementptr i64, ptr [[PTRB]], i64 8
+; CHECK-NEXT: [[PTR_ARM2:%.*]] = getelementptr i64, ptr [[PTRB]], i64 16
+; CHECK-NEXT: [[PTRC:%.*]] = select i1 [[COND2]], ptr [[PTR_ARM1]], ptr [[PTR_ARM2]]
+; CHECK-NEXT: store i64 [[VALUE]], ptr [[PTRC]], align 4
+; CHECK-NEXT: br label %[[MAINC]]
+; CHECK: [[MAINC]]:
+; CHECK-NEXT: store i64 [[VALUE]], ptr [[PTRD]], align 4
+; CHECK-NEXT: ret void
+;
+mainA:
+ %cond = icmp slt i64 %a, %b
+ br i1 %cond, label %ifA, label %mainB
+
+ifA:
+ %518 = load i64, ptr %ptrA
+ br label %mainB
+
+mainB:
+ %value = phi i64 [ %518, %ifA ], [ zeroinitializer, %mainA ]
+ br i1 %cond, label %ifB, label %mainC
+
+ifB:
+ %cond2 = icmp slt i64 %a, %c
+ br i1 %cond2, label %ifB_arm1, label %ifB_arm2
+
+ifB_arm1:
+ %ptr_arm1 = getelementptr i64, ptr %ptrB, i64 8
+ br label %ifB_join
+
+ifB_arm2:
+ %ptr_arm2 = getelementptr i64, ptr %ptrB, i64 16
+ br label %ifB_join
+
+ifB_join:
+ %ptrC = phi ptr [ %ptr_arm1, %ifB_arm1 ], [ %ptr_arm2, %ifB_arm2 ]
+ store i64 %value, ptr %ptrC
+ br label %mainC
+
+mainC:
+ store i64 %value, ptr %ptrD
+ ret void
+}
+
diff --git a/llvm/test/Transforms/SimplifyCFG/jump-threading-max-jump-threading-live-blocks.ll b/llvm/test/Transforms/SimplifyCFG/jump-threading-max-jump-threading-live-blocks.ll
new file mode 100644
index 0000000000000..686869348019d
--- /dev/null
+++ b/llvm/test/Transforms/SimplifyCFG/jump-threading-max-jump-threading-live-blocks.ll
@@ -0,0 +1,95 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 5
+; RUN: opt -passes=simplifycfg -S -max-jump-threading-live-blocks=3 < %s | FileCheck %s --check-prefixes=CHECK_LIMIT_3
+; RUN: opt -passes=simplifycfg -S -max-jump-threading-live-blocks=4 < %s | FileCheck %s --check-prefixes=CHECK_LIMIT_4
+
+; Test option -max-jump-threading-live-blocks=<num>
+
+define void @testB(ptr %ptrA, ptr %ptrB, i64 %a, i64 %b, i64 %c) {
+; CHECK_LIMIT_3-LABEL: define void @testB(
+; CHECK_LIMIT_3-SAME: ptr [[PTRA:%.*]], ptr [[PTRB:%.*]], i64 [[A:%.*]], i64 [[B:%.*]], i64 [[C:%.*]]) {
+; CHECK_LIMIT_3-NEXT: [[MAINA:.*]]:
+; CHECK_LIMIT_3-NEXT: [[COND:%.*]] = icmp slt i64 [[A]], [[B]]
+; CHECK_LIMIT_3-NEXT: br i1 [[COND]], label %[[IFA:.*]], label %[[MAINB:.*]]
+; CHECK_LIMIT_3: [[IFA]]:
+; CHECK_LIMIT_3-NEXT: [[TMP0:%.*]] = load i64, ptr [[PTRA]], align 4
+; CHECK_LIMIT_3-NEXT: br label %[[MAINB]]
+; CHECK_LIMIT_3: [[MAINB]]:
+; CHECK_LIMIT_3-NEXT: [[VALUE:%.*]] = phi i64 [ [[TMP0]], %[[IFA]] ], [ 0, %[[MAINA]] ]
+; CHECK_LIMIT_3-NEXT: br i1 [[COND]], label %[[IFB:.*]], label %[[MAINC:.*]]
+; CHECK_LIMIT_3: [[IFB]]:
+; CHECK_LIMIT_3-NEXT: [[COND2:%.*]] = icmp slt i64 [[A]], [[C]]
+; CHECK_LIMIT_3-NEXT: br i1 [[COND2]], label %[[IFB_ARM1:.*]], label %[[IFB_ARM2:.*]]
+; CHECK_LIMIT_3: [[IFB_ARM1]]:
+; CHECK_LIMIT_3-NEXT: [[PTR_ARM1:%.*]] = getelementptr i64, ptr [[PTRB]], i64 8
+; CHECK_LIMIT_3-NEXT: store i128 0, ptr [[PTR_ARM1]], align 4
+; CHECK_LIMIT_3-NEXT: br label %[[IFB_JOIN:.*]]
+; CHECK_LIMIT_3: [[IFB_ARM2]]:
+; CHECK_LIMIT_3-NEXT: [[PTR_ARM2:%.*]] = getelementptr i64, ptr [[PTRB]], i64 16
+; CHECK_LIMIT_3-NEXT: store i128 0, ptr [[PTR_ARM2]], align 4
+; CHECK_LIMIT_3-NEXT: br label %[[IFB_JOIN]]
+; CHECK_LIMIT_3: [[IFB_JOIN]]:
+; CHECK_LIMIT_3-NEXT: [[PTRC:%.*]] = phi ptr [ [[PTR_ARM1]], %[[IFB_ARM1]] ], [ [[PTR_ARM2]], %[[IFB_ARM2]] ]
+; CHECK_LIMIT_3-NEXT: store i64 [[VALUE]], ptr [[PTRC]], align 4
+; CHECK_LIMIT_3-NEXT: br label %[[MAINC]]
+; CHECK_LIMIT_3: [[MAINC]]:
+; CHECK_LIMIT_3-NEXT: ret void
+;
+; CHECK_LIMIT_4-LABEL: define void @testB(
+; CHECK_LIMIT_4-SAME: ptr [[PTRA:%.*]], ptr [[PTRB:%.*]], i64 [[A:%.*]], i64 [[B:%.*]], i64 [[C:%.*]]) {
+; CHECK_LIMIT_4-NEXT: [[MAINA:.*:]]
+; CHECK_LIMIT_4-NEXT: [[COND:%.*]] = icmp slt i64 [[A]], [[B]]
+; CHECK_LIMIT_4-NEXT: br i1 [[COND]], label %[[IFA:.*]], label %[[MAINC:.*]]
+; CHECK_LIMIT_4: [[IFA]]:
+; CHECK_LIMIT_4-NEXT: [[TMP0:%.*]] = load i64, ptr [[PTRA]], align 4
+; CHECK_LIMIT_4-NEXT: [[COND2:%.*]] = icmp slt i64 [[A]], [[C]]
+; CHECK_LIMIT_4-NEXT: br i1 [[COND2]], label %[[IFB_ARM1:.*]], label %[[IFB_ARM2:.*]]
+; CHECK_LIMIT_4: [[IFB_ARM1]]:
+; CHECK_LIMIT_4-NEXT: [[PTR_ARM1:%.*]] = getelementptr i64, ptr [[PTRB]], i64 8
+; CHECK_LIMIT_4-NEXT: store i128 0, ptr [[PTR_ARM1]], align 4
+; CHECK_LIMIT_4-NEXT: br label %[[IFB_JOIN:.*]]
+; CHECK_LIMIT_4: [[IFB_ARM2]]:
+; CHECK_LIMIT_4-NEXT: [[PTR_ARM2:%.*]] = getelementptr i64, ptr [[PTRB]], i64 16
+; CHECK_LIMIT_4-NEXT: store i128 0, ptr [[PTR_ARM2]], align 4
+; CHECK_LIMIT_4-NEXT: br label %[[IFB_JOIN]]
+; CHECK_LIMIT_4: [[IFB_JOIN]]:
+; CHECK_LIMIT_4-NEXT: [[PTRC:%.*]] = phi ptr [ [[PTR_ARM1]], %[[IFB_ARM1]] ], [ [[PTR_ARM2]], %[[IFB_ARM2]] ]
+; CHECK_LIMIT_4-NEXT: store i64 [[TMP0]], ptr [[PTRC]], align 4
+; CHECK_LIMIT_4-NEXT: br label %[[MAINC]]
+; CHECK_LIMIT_4: [[MAINC]]:
+; CHECK_LIMIT_4-NEXT: ret void
+;
+mainA:
+ %cond = icmp slt i64 %a, %b
+ br i1 %cond, label %ifA, label %mainB
+
+ifA:
+ %518 = load i64, ptr %ptrA
+ br label %mainB
+
+; Use of %value is not in either immediate destination of mainB.
+mainB:
+ %value = phi i64 [ %518, %ifA ], [ zeroinitializer, %mainA ]
+ br i1 %cond, label %ifB, label %mainC
+
+ifB:
+ %cond2 = icmp slt i64 %a, %c
+ br i1 %cond2, label %ifB_arm1, label %ifB_arm2
+
+ifB_arm1:
+ %ptr_arm1 = getelementptr i64, ptr %ptrB, i64 8
+ store i128 0, ptr %ptr_arm1
+ br label %ifB_join
+
+ifB_arm2:
+ %ptr_arm2 = getelementptr i64, ptr %ptrB, i64 16
+ store i128 0, ptr %ptr_arm2
+ br label %ifB_join
+
+ifB_join:
+ %ptrC = phi ptr [ %ptr_arm1, %ifB_arm1 ], [ %ptr_arm2, %ifB_arm2 ]
+ store i64 %value, ptr %ptrC
+ br label %mainC
+
+mainC:
+ ret void
+}
>From aca42203f37fc6b00be23ade2e1d99db8ab36cea Mon Sep 17 00:00:00 2001
From: John Lu <John.Lu at amd.com>
Date: Fri, 22 May 2026 14:03:33 -0500
Subject: [PATCH 2/2] Add testcase showing SROA increasing phi-node count
Signed-off-by: John Lu <John.Lu at amd.com>
---
llvm/test/CodeGen/AMDGPU/sroa-phi-nodes.ll | 167 +++++++++++++++++++++
1 file changed, 167 insertions(+)
create mode 100644 llvm/test/CodeGen/AMDGPU/sroa-phi-nodes.ll
diff --git a/llvm/test/CodeGen/AMDGPU/sroa-phi-nodes.ll b/llvm/test/CodeGen/AMDGPU/sroa-phi-nodes.ll
new file mode 100644
index 0000000000000..3ca0c4cad6ade
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/sroa-phi-nodes.ll
@@ -0,0 +1,167 @@
+; RUN: opt -mtriple=amdgcn -O2 -S %s -o - -stop-after=sroa | FileCheck %s
+
+; This testcase is dervied from warp's reduce.cu.
+;
+; Pathological testcase where SROA increases the number of phi-nodes from ~250 to 32K.
+; This increase in phi-nodes led to a large increase in compile time in SSAUpdater
+; when called by StructurizeCFG when compiling reduce.cu. See large-phi-search.ll for a
+; minimized testcase for the SSAUpdater slowdown.
+;
+; Note that this large increase in phi count by SROA depends on allowing jump-threading
+; through blocks that define values that are live outside of the block.
+; There is no large increase in phi count with:
+;
+; opt -mtriple=amdgcn -O2 -S %s -o - -stop-after=sroa -max-jump-threading-live-blocks=0
+
+
+; CHECK-LABEL: @_func7(
+
+%"cub_inner_product_iterator" = type <{ ptr, ptr, i32, i32, i32, [4 x i8] }>
+
+; Function Attrs: convergent inlinehint mustprogress nounwind
+define noundef double @_func(ptr noundef nonnull align 8 dereferenceable(8) %var0, ptr noundef nonnull align 8 dereferenceable(8) %var1) #0 {
+ %var3 = load double, ptr %var0, align 8
+ %var4 = load double, ptr %var1, align 8
+ %var5 = fadd contract double %var3, %var4
+ ret double %var5
+}
+
+; Function Attrs: convergent inlinehint mustprogress nounwind
+define noundef double @_func3(ptr noundef nonnull align 8 dereferenceable(8) %var0, ptr noundef nonnull align 8 dereferenceable(8) %var1) #0 {
+ %var3 = call contract noundef double @_func(ptr noundef nonnull align 8 dereferenceable(8) %var0, ptr noundef nonnull align 8 dereferenceable(8) %var1) #3
+ ret double %var3
+}
+
+; Function Attrs: convergent inlinehint mustprogress nounwind
+declare void @_func2(double noundef) #0
+
+; Function Attrs: alwaysinline convergent mustprogress nounwind
+define void @_func4(ptr noundef nonnull align 8 dereferenceable(2048) %var0) #1 {
+ call void @_func5(ptr noundef nonnull align 8 dereferenceable(2048) %var0) #3
+ ret void
+}
+
+; Function Attrs: alwaysinline convergent mustprogress nounwind
+define void @_func5(ptr noundef nonnull align 8 dereferenceable(2048) %var0) #1 {
+ call void @_func6(ptr noundef nonnull align 8 dereferenceable(2048) %var0) #3
+ ret void
+}
+
+; Function Attrs: convergent inlinehint mustprogress nounwind
+define void @_func6(ptr noundef nonnull align 8 dereferenceable(2048) %var0) #0 {
+block1:
+ %var2 = alloca double, align 8, addrspace(5)
+ %var3 = addrspacecast ptr addrspace(5) %var2 to ptr
+ store double 0.000000e+00, ptr %var3, align 8
+ br label %block4
+
+block4: ; preds = %block8, %block1
+ %var5 = phi double [ 0.000000e+00, %block1 ], [ %var11, %block8 ]
+ %.0 = phi i32 [ 0, %block1 ], [ %var12, %block8 ]
+ %var6 = icmp samesign ult i32 %.0, 256
+ br i1 %var6, label %block8, label %block7
+
+block7: ; preds = %block4
+ call void @_func2(double noundef %var5) #3
+ ret void
+
+block8: ; preds = %block4
+ %var9 = zext nneg i32 %.0 to i64
+ %var10 = getelementptr inbounds nuw [8 x i8], ptr %var0, i64 %var9
+ %var11 = call contract noundef double @_func3(ptr noundef nonnull align 8 dereferenceable(8) %var3, ptr noundef nonnull align 8 dereferenceable(8) %var10) #3
+ store double %var11, ptr %var3, align 8
+ %var12 = add nuw nsw i32 %.0, 1
+ br label %block4, !llvm.loop !0
+}
+
+; Function Attrs: convergent mustprogress nounwind
+declare %"cub_inner_product_iterator" @_func8() #2
+
+; Function Attrs: convergent mustprogress nounwind
+define noundef double @_func9(ptr noundef nonnull align 8 dereferenceable(28) %var0, i64 noundef %var1) #2 {
+ %var3 = call contract noundef double @_func10(ptr noundef nonnull align 8 dereferenceable(28) %var0, i64 noundef %var1) #3
+ ret double %var3
+}
+
+; Function Attrs: convergent mustprogress nounwind
+define noundef double @_func10(ptr noundef nonnull align 8 dereferenceable(28) %var0, i64 noundef %var1) #2 {
+block2:
+ %var3 = getelementptr inbounds nuw i8, ptr %var0, i64 16
+ %var4 = load i32, ptr %var3, align 8
+ %var5 = sext i32 %var4 to i64
+ %var6 = mul nsw i64 %var1, %var5
+ %var7 = getelementptr inbounds [8 x i8], ptr null, i64 %var6
+ br label %block8
+
+block8: ; preds = %block13, %block2
+ %.05 = phi double [ 0.000000e+00, %block2 ], [ 1.000000e+00, %block13 ]
+ %.0 = phi i32 [ 0, %block2 ], [ %var16, %block13 ]
+ %var9 = getelementptr inbounds nuw i8, ptr %var0, i64 24
+ %var10 = load i32, ptr %var9, align 8
+ %var11 = icmp slt i32 %.0, %var10
+ br i1 %var11, label %block13, label %block12
+
+block12: ; preds = %block8
+ ret double %.05
+
+block13: ; preds = %block8
+ %var14 = call contract noundef double @_func11(ptr noundef nonnull align 8 dereferenceable(8) %var7) #3
+ %var15 = fadd contract double 1.000000e+00, 0.000000e+00
+ %var16 = add nuw nsw i32 1, 1
+ br label %block8
+}
+
+; Function Attrs: convergent mustprogress nounwind
+declare noundef double @_func11(ptr noundef nonnull align 8 dereferenceable(8)) #2
+
+; Function Attrs: alwaysinline convergent mustprogress nounwind
+define void @_func7() #1 {
+ %var1 = alloca [256 x double], align 16, addrspace(5)
+ %var2 = addrspacecast ptr addrspace(5) %var1 to ptr
+ call void @_func12(ptr noundef nonnull align 8 dereferenceable(2048) %var2) #3
+ call void @_func4(ptr noundef nonnull align 8 dereferenceable(2048) %var2) #3
+ ret void
+}
+
+; Function Attrs: convergent inlinehint mustprogress nounwind
+define void @_func12(ptr noundef nonnull align 8 dereferenceable(2048) %var0) #0 {
+block1:
+ %var2 = alloca %"cub_inner_product_iterator", align 8, addrspace(5)
+ %var3 = addrspacecast ptr addrspace(5) %var2 to ptr
+ %var4 = call %"cub_inner_product_iterator" @_func8() #3
+ %.fca.2.extract = extractvalue %"cub_inner_product_iterator" %var4, 2
+ %.fca.4.extract = extractvalue %"cub_inner_product_iterator" %var4, 4
+ %.sroa.3.0..sroa_idx = getelementptr inbounds nuw i8, ptr %var3, i64 16
+ store i32 %.fca.2.extract, ptr %.sroa.3.0..sroa_idx, align 8
+ %.sroa.5.0..sroa_idx = getelementptr inbounds nuw i8, ptr %var3, i64 24
+ store i32 %.fca.4.extract, ptr %.sroa.5.0..sroa_idx, align 8
+ br label %block5
+
+block5: ; preds = %block8, %block1
+ %.0 = phi i32 [ 0, %block1 ], [ %var14, %block8 ]
+ %var6 = icmp samesign ult i32 %.0, 256
+ br i1 %var6, label %block8, label %block7
+
+block7: ; preds = %block5
+ ret void
+
+block8: ; preds = %block5
+ %var9 = shl nuw nsw i32 %.0, 0
+ %var10 = zext nneg i32 %var9 to i64
+ %var11 = call contract noundef double @_func9(ptr noundef nonnull align 8 dereferenceable(28) %var3, i64 noundef %var10) #3
+ %var12 = zext nneg i32 %.0 to i64
+ %var13 = getelementptr inbounds nuw [8 x i8], ptr %var0, i64 %var12
+ store double %var11, ptr %var13, align 8
+ %var14 = add nuw nsw i32 %.0, 1
+ br label %block5, !llvm.loop !3
+}
+
+attributes #0 = { convergent inlinehint mustprogress nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx942" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+lerp-inst,+mai-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+wavefrontsize64,+xf32-insts" }
+attributes #1 = { alwaysinline convergent mustprogress nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx942" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+lerp-inst,+mai-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+wavefrontsize64,+xf32-insts" }
+attributes #2 = { convergent mustprogress nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx942" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+lerp-inst,+mai-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+wavefrontsize64,+xf32-insts" }
+attributes #3 = { convergent nounwind }
+
+!0 = distinct !{!0, !1, !2}
+!1 = !{!"llvm.loop.mustprogress"}
+!2 = !{!"llvm.loop.unroll.enable"}
+!3 = distinct !{!3, !1, !2}
More information about the cfe-commits
mailing list