[clang] [llvm] [SimplifyCFG] Hoist common code for multi-case destinations of switch (PR #165700)
Kunqiu Chen via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 30 05:54:06 PDT 2025
https://github.com/Camsyn updated https://github.com/llvm/llvm-project/pull/165700
>From 091cc77fb2993a9e3df4c3583da4a92ac9d5d166 Mon Sep 17 00:00:00 2001
From: Camsyn <camsyn at foxmail.com>
Date: Thu, 30 Oct 2025 16:49:01 +0800
Subject: [PATCH 1/4] Add a test before commit
---
.../SimplifyCFG/hoist-common-code.ll | 47 +++++++++++++++++++
1 file changed, 47 insertions(+)
diff --git a/llvm/test/Transforms/SimplifyCFG/hoist-common-code.ll b/llvm/test/Transforms/SimplifyCFG/hoist-common-code.ll
index 98c0599ab209c..e70d6fb7d8bb2 100644
--- a/llvm/test/Transforms/SimplifyCFG/hoist-common-code.ll
+++ b/llvm/test/Transforms/SimplifyCFG/hoist-common-code.ll
@@ -602,3 +602,50 @@ bar:
call void @bar()
ret void
}
+
+define void @test_switch_with_multicases_dest(i64 %i, ptr %Q) {
+; CHECK-LABEL: @test_switch_with_multicases_dest(
+; CHECK-NEXT: switch i64 [[I:%.*]], label [[BB0:%.*]] [
+; CHECK-NEXT: i64 1, label [[BB1:%.*]]
+; CHECK-NEXT: i64 2, label [[BB2:%.*]]
+; CHECK-NEXT: i64 3, label [[BB2]]
+; CHECK-NEXT: ]
+; CHECK: common.ret:
+; CHECK-NEXT: ret void
+; CHECK: bb0:
+; CHECK-NEXT: store i32 1, ptr [[Q:%.*]], align 4
+; CHECK-NEXT: [[A:%.*]] = load i32, ptr [[Q]], align 4
+; CHECK-NEXT: call void @bar(i32 [[A]])
+; CHECK-NEXT: br label [[COMMON_RET:%.*]]
+; CHECK: bb1:
+; CHECK-NEXT: store i32 1, ptr [[Q]], align 4
+; CHECK-NEXT: [[B:%.*]] = load i32, ptr [[Q]], align 4
+; CHECK-NEXT: call void @bar(i32 [[B]])
+; CHECK-NEXT: br label [[COMMON_RET]]
+; CHECK: bb2:
+; CHECK-NEXT: store i32 1, ptr [[Q]], align 4
+; CHECK-NEXT: [[C:%.*]] = load i32, ptr [[Q]], align 4
+; CHECK-NEXT: call void @bar(i32 [[C]])
+; CHECK-NEXT: br label [[COMMON_RET]]
+;
+ switch i64 %i, label %bb0 [
+ i64 1, label %bb1
+ i64 2, label %bb2
+ i64 3, label %bb2
+ ]
+bb0: ; preds = %0
+ store i32 1, ptr %Q
+ %A = load i32, ptr %Q ; <i32> [#uses=1]
+ call void @bar( i32 %A )
+ ret void
+bb1: ; preds = %0
+ store i32 1, ptr %Q
+ %B = load i32, ptr %Q ; <i32> [#uses=1]
+ call void @bar( i32 %B )
+ ret void
+bb2: ; preds = %0
+ store i32 1, ptr %Q
+ %C = load i32, ptr %Q ; <i32> [#uses=1]
+ call void @bar( i32 %C )
+ ret void
+}
>From 9877726e958daffa8908f292fc7d3ebce853dd29 Mon Sep 17 00:00:00 2001
From: Camsyn <camsyn at foxmail.com>
Date: Thu, 30 Oct 2025 19:32:45 +0800
Subject: [PATCH 2/4] [SimplifyCFG] Hoist common code for multi-cases dest
---
llvm/lib/Transforms/Utils/SimplifyCFG.cpp | 30 ++++++++++++++++-------
1 file changed, 21 insertions(+), 9 deletions(-)
diff --git a/llvm/lib/Transforms/Utils/SimplifyCFG.cpp b/llvm/lib/Transforms/Utils/SimplifyCFG.cpp
index b03fb6213d61c..4d2a840bf059f 100644
--- a/llvm/lib/Transforms/Utils/SimplifyCFG.cpp
+++ b/llvm/lib/Transforms/Utils/SimplifyCFG.cpp
@@ -1866,10 +1866,18 @@ bool SimplifyCFGOpt::hoistCommonCodeFromSuccessors(Instruction *TI,
// If either of the blocks has it's address taken, then we can't do this fold,
// because the code we'd hoist would no longer run when we jump into the block
// by it's address.
+ SmallVector<BasicBlock *, 4> UniqSuccessors;
+ UniqSuccessors.reserve(BB->getTerminator()->getNumSuccessors());
for (auto *Succ : successors(BB)) {
if (Succ->hasAddressTaken())
return false;
- if (Succ->getSinglePredecessor())
+ // Collect all unique successors, using vec instead of set to preserve
+ // order.
+ if (find(UniqSuccessors, Succ) == UniqSuccessors.end())
+ UniqSuccessors.push_back(Succ);
+ // Use getUniquePredecessor instead of getSinglePredecessor to support
+ // multi-cases successors in switch.
+ if (Succ->getUniquePredecessor())
continue;
// If Succ has >1 predecessors, continue to check if the Succ contains only
// one `unreachable` inst. Since executing `unreachable` inst is an UB, we
@@ -1882,7 +1890,7 @@ bool SimplifyCFGOpt::hoistCommonCodeFromSuccessors(Instruction *TI,
// The second of pair is a SkipFlags bitmask.
using SuccIterPair = std::pair<BasicBlock::iterator, unsigned>;
SmallVector<SuccIterPair, 8> SuccIterPairs;
- for (auto *Succ : successors(BB)) {
+ for (auto *Succ : UniqSuccessors) {
BasicBlock::iterator SuccItr = Succ->begin();
if (isa<PHINode>(*SuccItr))
return false;
@@ -1893,19 +1901,19 @@ bool SimplifyCFGOpt::hoistCommonCodeFromSuccessors(Instruction *TI,
// Check if all instructions in the successor blocks match. This allows
// hoisting all instructions and removing the blocks we are hoisting from,
// so does not add any new instructions.
- SmallVector<BasicBlock *> Succs = to_vector(successors(BB));
+
// Check if sizes and terminators of all successors match.
- bool AllSame = none_of(Succs, [&Succs](BasicBlock *Succ) {
- Instruction *Term0 = Succs[0]->getTerminator();
+ bool AllSame = none_of(UniqSuccessors, [&UniqSuccessors](BasicBlock *Succ) {
+ Instruction *Term0 = UniqSuccessors[0]->getTerminator();
Instruction *Term = Succ->getTerminator();
return !Term->isSameOperationAs(Term0) ||
!equal(Term->operands(), Term0->operands()) ||
- Succs[0]->size() != Succ->size();
+ UniqSuccessors[0]->size() != Succ->size();
});
if (!AllSame)
return false;
if (AllSame) {
- LockstepReverseIterator<true> LRI(Succs);
+ LockstepReverseIterator<true> LRI(UniqSuccessors);
while (LRI.isValid()) {
Instruction *I0 = (*LRI)[0];
if (any_of(*LRI, [I0](Instruction *I) {
@@ -2156,9 +2164,13 @@ bool SimplifyCFGOpt::hoistSuccIdenticalTerminatorToSwitchOrIf(
Updates.push_back({DominatorTree::Insert, TIParent, Succ});
}
- if (DTU)
- for (BasicBlock *Succ : successors(TI))
+ if (DTU) {
+ // TI might be a switch with multi-cases destination, so we need to care for
+ // the duplication of successors.
+ SmallPtrSet<BasicBlock *, 4> UniqSuccs(llvm::from_range, successors(TI));
+ for (BasicBlock *Succ : UniqSuccs)
Updates.push_back({DominatorTree::Delete, TIParent, Succ});
+ }
eraseTerminatorAndDCECond(TI);
if (DTU)
>From 93781e5daa2ad1474182be6513a24dc588f2106d Mon Sep 17 00:00:00 2001
From: Camsyn <camsyn at foxmail.com>
Date: Thu, 30 Oct 2025 19:33:28 +0800
Subject: [PATCH 3/4] Regenerate test after commit
---
.../SimplifyCFG/hoist-common-code.ll | 21 ++-----------------
1 file changed, 2 insertions(+), 19 deletions(-)
diff --git a/llvm/test/Transforms/SimplifyCFG/hoist-common-code.ll b/llvm/test/Transforms/SimplifyCFG/hoist-common-code.ll
index e70d6fb7d8bb2..68103e874923e 100644
--- a/llvm/test/Transforms/SimplifyCFG/hoist-common-code.ll
+++ b/llvm/test/Transforms/SimplifyCFG/hoist-common-code.ll
@@ -605,28 +605,11 @@ bar:
define void @test_switch_with_multicases_dest(i64 %i, ptr %Q) {
; CHECK-LABEL: @test_switch_with_multicases_dest(
-; CHECK-NEXT: switch i64 [[I:%.*]], label [[BB0:%.*]] [
-; CHECK-NEXT: i64 1, label [[BB1:%.*]]
-; CHECK-NEXT: i64 2, label [[BB2:%.*]]
-; CHECK-NEXT: i64 3, label [[BB2]]
-; CHECK-NEXT: ]
-; CHECK: common.ret:
-; CHECK-NEXT: ret void
-; CHECK: bb0:
+; CHECK-NEXT: common.ret:
; CHECK-NEXT: store i32 1, ptr [[Q:%.*]], align 4
-; CHECK-NEXT: [[A:%.*]] = load i32, ptr [[Q]], align 4
-; CHECK-NEXT: call void @bar(i32 [[A]])
-; CHECK-NEXT: br label [[COMMON_RET:%.*]]
-; CHECK: bb1:
-; CHECK-NEXT: store i32 1, ptr [[Q]], align 4
-; CHECK-NEXT: [[B:%.*]] = load i32, ptr [[Q]], align 4
-; CHECK-NEXT: call void @bar(i32 [[B]])
-; CHECK-NEXT: br label [[COMMON_RET]]
-; CHECK: bb2:
-; CHECK-NEXT: store i32 1, ptr [[Q]], align 4
; CHECK-NEXT: [[C:%.*]] = load i32, ptr [[Q]], align 4
; CHECK-NEXT: call void @bar(i32 [[C]])
-; CHECK-NEXT: br label [[COMMON_RET]]
+; CHECK-NEXT: ret void
;
switch i64 %i, label %bb0 [
i64 1, label %bb1
>From c73ac96e49960bb1d3d100ab44b4c4555d8ca88e Mon Sep 17 00:00:00 2001
From: Camsyn <camsyn at foxmail.com>
Date: Thu, 30 Oct 2025 20:53:31 +0800
Subject: [PATCH 4/4] Regenerate another test
---
clang/test/Headers/__clang_hip_math.hip | 381 +++++++++++-------------
1 file changed, 180 insertions(+), 201 deletions(-)
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index 7e2691633c215..318e60a427aee 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -244,74 +244,71 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
// CHECK: [[IF_THEN_I]]:
// CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds nuw i8, ptr [[P]], i64 1
// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[INCDEC_PTR_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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_I_I_PREHEADER]]:
-// CHECK-NEXT: [[TMP3:%.*]] = load i8, ptr [[INCDEC_PTR_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// CHECK-NEXT: [[CMP_NOT_I_I14:%.*]] = icmp eq i8 [[TMP3]], 0
// CHECK-NEXT: br i1 [[CMP_NOT_I_I14]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], label %[[WHILE_BODY_I_I:.*]]
// CHECK: [[IF_THEN5_I]]:
-// CHECK-NEXT: [[TMP4:%.*]] = load i8, ptr [[INCDEC_PTR_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// CHECK-NEXT: [[CMP_NOT_I30_I9:%.*]] = icmp eq i8 [[TMP4]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I30_I9]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], label %[[WHILE_BODY_I31_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: [[TMP5:%.*]] = phi i8 [ [[TMP9:%.*]], %[[IF_END31_I_I:.*]] ], [ [[TMP4]], %[[IF_THEN5_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: [[TMP6:%.*]] = add i8 [[TMP5]], -48
-// CHECK-NEXT: [[OR_COND_I32_I:%.*]] = icmp ult i8 [[TMP6]], 10
+// 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: [[TMP7:%.*]] = add i8 [[TMP5]], -97
-// CHECK-NEXT: [[OR_COND33_I_I:%.*]] = icmp ult i8 [[TMP7]], 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: [[TMP8:%.*]] = add i8 [[TMP5]], -65
-// CHECK-NEXT: [[OR_COND34_I_I:%.*]] = icmp ult i8 [[TMP8]], 6
+// 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 [ -48, %[[WHILE_BODY_I31_I]] ], [ -87, %[[IF_ELSE_I_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 [[TMP5]] to i64
+// 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_I34_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I28_I10]], i64 1
-// CHECK-NEXT: [[TMP9]] = load i8, ptr [[INCDEC_PTR_I34_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// CHECK-NEXT: [[CMP_NOT_I30_I:%.*]] = icmp eq i8 [[TMP9]], 0
+// CHECK-NEXT: [[TMP7]] = load i8, ptr [[INCDEC_PTR_I34_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP13]]
// CHECK: [[WHILE_BODY_I_I]]:
-// CHECK-NEXT: [[TMP10:%.*]] = phi i8 [ [[TMP12:%.*]], %[[IF_THEN_I_I:.*]] ], [ [[TMP3]], %[[WHILE_COND_I_I_PREHEADER]] ]
+// 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: [[TMP11:%.*]] = and i8 [[TMP10]], -8
-// CHECK-NEXT: [[OR_COND_I_I:%.*]] = icmp eq i8 [[TMP11]], 48
+// 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_I16]], 3
-// CHECK-NEXT: [[CONV5_I_I:%.*]] = zext nneg i8 [[TMP10]] to i64
+// 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_I15]], i64 1
-// CHECK-NEXT: [[TMP12]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// CHECK-NEXT: [[CMP_NOT_I_I:%.*]] = icmp eq i8 [[TMP12]], 0
+// CHECK-NEXT: [[TMP10]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP9]]
// CHECK: [[WHILE_BODY_I18_I]]:
-// CHECK-NEXT: [[TMP13:%.*]] = phi i8 [ [[TMP15:%.*]], %[[IF_THEN_I21_I:.*]] ], [ [[TMP1]], %[[WHILE_COND_I14_I_PREHEADER]] ]
+// 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: [[TMP14:%.*]] = add i8 [[TMP13]], -48
-// CHECK-NEXT: [[OR_COND_I19_I:%.*]] = icmp ult i8 [[TMP14]], 10
+// 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 [[TMP13]] to i64
+// 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: [[TMP15]] = load i8, ptr [[INCDEC_PTR_I26_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// CHECK-NEXT: [[CMP_NOT_I17_I:%.*]] = icmp eq i8 [[TMP15]], 0
+// CHECK-NEXT: [[TMP13]] = load i8, ptr [[INCDEC_PTR_I26_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP12]]
// CHECK: [[_ZL15__MAKE_MANTISSAPKC_EXIT]]:
// CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, %[[WHILE_COND_I_I_PREHEADER]] ], [ 0, %[[IF_THEN5_I]] ], [ 0, %[[WHILE_COND_I14_I_PREHEADER]] ], [ [[SUB_I_I]], %[[IF_THEN_I_I]] ], [ 0, %[[WHILE_BODY_I_I]] ], [ [[ADD28_I_I]], %[[IF_END31_I_I]] ], [ 0, %[[IF_ELSE17_I_I]] ], [ [[SUB_I25_I]], %[[IF_THEN_I21_I]] ], [ 0, %[[WHILE_BODY_I18_I]] ]
@@ -4265,82 +4262,79 @@ extern "C" __device__ double test_modf(double x, double* y) {
// DEFAULT: [[IF_THEN_I_I]]:
// DEFAULT-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[TAG]], i64 1
// DEFAULT-NEXT: [[TMP2:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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_I_I_I_PREHEADER]]:
-// DEFAULT-NEXT: [[TMP3:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// DEFAULT-NEXT: [[CMP_NOT_I_I_I14:%.*]] = icmp eq i8 [[TMP3]], 0
// 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: [[TMP4:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// DEFAULT-NEXT: [[CMP_NOT_I30_I_I9:%.*]] = icmp eq i8 [[TMP4]], 0
-// DEFAULT-NEXT: br i1 [[CMP_NOT_I30_I_I9]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I31_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: [[TMP5:%.*]] = phi i8 [ [[TMP9:%.*]], %[[IF_END31_I_I_I:.*]] ], [ [[TMP4]], %[[IF_THEN5_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: [[TMP6:%.*]] = add i8 [[TMP5]], -48
-// DEFAULT-NEXT: [[OR_COND_I32_I_I:%.*]] = icmp ult i8 [[TMP6]], 10
+// 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: [[TMP7:%.*]] = add i8 [[TMP5]], -97
-// DEFAULT-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP7]], 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: [[TMP8:%.*]] = add i8 [[TMP5]], -65
-// DEFAULT-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP8]], 6
+// 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 [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -87, %[[IF_ELSE_I_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 [[TMP5]] to i64
+// 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_I34_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I28_I_I10]], i64 1
-// DEFAULT-NEXT: [[TMP9]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// DEFAULT-NEXT: [[CMP_NOT_I30_I_I:%.*]] = icmp eq i8 [[TMP9]], 0
+// DEFAULT-NEXT: [[TMP7]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP13]]
// DEFAULT: [[WHILE_BODY_I_I_I]]:
-// DEFAULT-NEXT: [[TMP10:%.*]] = phi i8 [ [[TMP12:%.*]], %[[IF_THEN_I_I_I:.*]] ], [ [[TMP3]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// 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: [[TMP11:%.*]] = and i8 [[TMP10]], -8
-// DEFAULT-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP11]], 48
+// 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_I16]], 3
-// DEFAULT-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP10]] to i64
+// 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_I15]], i64 1
-// DEFAULT-NEXT: [[TMP12]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// DEFAULT-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP12]], 0
+// DEFAULT-NEXT: [[TMP10]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP9]]
// DEFAULT: [[WHILE_BODY_I18_I_I]]:
-// DEFAULT-NEXT: [[TMP13:%.*]] = phi i8 [ [[TMP15:%.*]], %[[IF_THEN_I21_I_I:.*]] ], [ [[TMP1]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// 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: [[TMP14:%.*]] = add i8 [[TMP13]], -48
-// DEFAULT-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP14]], 10
+// 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 [[TMP13]] to i64
+// 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: [[TMP15]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// DEFAULT-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP15]], 0
+// DEFAULT-NEXT: [[TMP13]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP12]]
// DEFAULT: [[_ZL4NANFPKC_EXIT]]:
// DEFAULT-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_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: [[TMP16:%.*]] = bitcast i32 [[BF_SET9_I]] to float
-// DEFAULT-NEXT: ret float [[TMP16]]
+// 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]] {
@@ -4360,82 +4354,79 @@ extern "C" __device__ double test_modf(double x, double* y) {
// APPROX: [[IF_THEN_I_I]]:
// APPROX-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[TAG]], i64 1
// APPROX-NEXT: [[TMP2:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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_I_I_I_PREHEADER]]:
-// APPROX-NEXT: [[TMP3:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// APPROX-NEXT: [[CMP_NOT_I_I_I14:%.*]] = icmp eq i8 [[TMP3]], 0
// 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: [[TMP4:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// APPROX-NEXT: [[CMP_NOT_I30_I_I9:%.*]] = icmp eq i8 [[TMP4]], 0
-// APPROX-NEXT: br i1 [[CMP_NOT_I30_I_I9]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I31_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: [[TMP5:%.*]] = phi i8 [ [[TMP9:%.*]], %[[IF_END31_I_I_I:.*]] ], [ [[TMP4]], %[[IF_THEN5_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: [[TMP6:%.*]] = add i8 [[TMP5]], -48
-// APPROX-NEXT: [[OR_COND_I32_I_I:%.*]] = icmp ult i8 [[TMP6]], 10
+// 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: [[TMP7:%.*]] = add i8 [[TMP5]], -97
-// APPROX-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP7]], 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: [[TMP8:%.*]] = add i8 [[TMP5]], -65
-// APPROX-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP8]], 6
+// 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 [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -87, %[[IF_ELSE_I_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 [[TMP5]] to i64
+// 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_I34_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I28_I_I10]], i64 1
-// APPROX-NEXT: [[TMP9]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// APPROX-NEXT: [[CMP_NOT_I30_I_I:%.*]] = icmp eq i8 [[TMP9]], 0
+// APPROX-NEXT: [[TMP7]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP13]]
// APPROX: [[WHILE_BODY_I_I_I]]:
-// APPROX-NEXT: [[TMP10:%.*]] = phi i8 [ [[TMP12:%.*]], %[[IF_THEN_I_I_I:.*]] ], [ [[TMP3]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// 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: [[TMP11:%.*]] = and i8 [[TMP10]], -8
-// APPROX-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP11]], 48
+// 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_I16]], 3
-// APPROX-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP10]] to i64
+// 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_I15]], i64 1
-// APPROX-NEXT: [[TMP12]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// APPROX-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP12]], 0
+// APPROX-NEXT: [[TMP10]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP9]]
// APPROX: [[WHILE_BODY_I18_I_I]]:
-// APPROX-NEXT: [[TMP13:%.*]] = phi i8 [ [[TMP15:%.*]], %[[IF_THEN_I21_I_I:.*]] ], [ [[TMP1]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// 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: [[TMP14:%.*]] = add i8 [[TMP13]], -48
-// APPROX-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP14]], 10
+// 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 [[TMP13]] to i64
+// 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: [[TMP15]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// APPROX-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP15]], 0
+// APPROX-NEXT: [[TMP13]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP12]]
// APPROX: [[_ZL4NANFPKC_EXIT]]:
// APPROX-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_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: [[TMP16:%.*]] = bitcast i32 [[BF_SET9_I]] to float
-// APPROX-NEXT: ret float [[TMP16]]
+// 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]] {
@@ -4450,82 +4441,79 @@ extern "C" __device__ double test_modf(double x, double* y) {
// NCRDIV: [[IF_THEN_I_I]]:
// NCRDIV-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[TAG]], i64 1
// NCRDIV-NEXT: [[TMP2:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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_I_I_I_PREHEADER]]:
-// NCRDIV-NEXT: [[TMP3:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// NCRDIV-NEXT: [[CMP_NOT_I_I_I14:%.*]] = icmp eq i8 [[TMP3]], 0
// 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: [[TMP4:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// NCRDIV-NEXT: [[CMP_NOT_I30_I_I9:%.*]] = icmp eq i8 [[TMP4]], 0
-// NCRDIV-NEXT: br i1 [[CMP_NOT_I30_I_I9]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I31_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: [[TMP5:%.*]] = phi i8 [ [[TMP9:%.*]], %[[IF_END31_I_I_I:.*]] ], [ [[TMP4]], %[[IF_THEN5_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: [[TMP6:%.*]] = add i8 [[TMP5]], -48
-// NCRDIV-NEXT: [[OR_COND_I32_I_I:%.*]] = icmp ult i8 [[TMP6]], 10
+// 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: [[TMP7:%.*]] = add i8 [[TMP5]], -97
-// NCRDIV-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP7]], 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: [[TMP8:%.*]] = add i8 [[TMP5]], -65
-// NCRDIV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP8]], 6
+// 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 [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -87, %[[IF_ELSE_I_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 [[TMP5]] to i64
+// 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_I34_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I28_I_I10]], i64 1
-// NCRDIV-NEXT: [[TMP9]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// NCRDIV-NEXT: [[CMP_NOT_I30_I_I:%.*]] = icmp eq i8 [[TMP9]], 0
+// NCRDIV-NEXT: [[TMP7]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP13]]
// NCRDIV: [[WHILE_BODY_I_I_I]]:
-// NCRDIV-NEXT: [[TMP10:%.*]] = phi i8 [ [[TMP12:%.*]], %[[IF_THEN_I_I_I:.*]] ], [ [[TMP3]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// 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: [[TMP11:%.*]] = and i8 [[TMP10]], -8
-// NCRDIV-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP11]], 48
+// 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_I16]], 3
-// NCRDIV-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP10]] to i64
+// 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_I15]], i64 1
-// NCRDIV-NEXT: [[TMP12]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// NCRDIV-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP12]], 0
+// NCRDIV-NEXT: [[TMP10]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP9]]
// NCRDIV: [[WHILE_BODY_I18_I_I]]:
-// NCRDIV-NEXT: [[TMP13:%.*]] = phi i8 [ [[TMP15:%.*]], %[[IF_THEN_I21_I_I:.*]] ], [ [[TMP1]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// 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: [[TMP14:%.*]] = add i8 [[TMP13]], -48
-// NCRDIV-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP14]], 10
+// 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 [[TMP13]] to i64
+// 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: [[TMP15]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// NCRDIV-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP15]], 0
+// NCRDIV-NEXT: [[TMP13]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP12]]
// NCRDIV: [[_ZL4NANFPKC_EXIT]]:
// NCRDIV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_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: [[TMP16:%.*]] = bitcast i32 [[BF_SET9_I]] to float
-// NCRDIV-NEXT: ret float [[TMP16]]
+// 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]] {
@@ -4628,81 +4616,78 @@ extern "C" __device__ float test_nanf(const char *tag) {
// DEFAULT: [[IF_THEN_I_I]]:
// DEFAULT-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[TAG]], i64 1
// DEFAULT-NEXT: [[TMP2:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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_I_I_I_PREHEADER]]:
-// DEFAULT-NEXT: [[TMP3:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// DEFAULT-NEXT: [[CMP_NOT_I_I_I14:%.*]] = icmp eq i8 [[TMP3]], 0
// 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: [[TMP4:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// DEFAULT-NEXT: [[CMP_NOT_I30_I_I9:%.*]] = icmp eq i8 [[TMP4]], 0
-// DEFAULT-NEXT: br i1 [[CMP_NOT_I30_I_I9]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I31_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: [[TMP5:%.*]] = phi i8 [ [[TMP9:%.*]], %[[IF_END31_I_I_I:.*]] ], [ [[TMP4]], %[[IF_THEN5_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: [[TMP6:%.*]] = add i8 [[TMP5]], -48
-// DEFAULT-NEXT: [[OR_COND_I32_I_I:%.*]] = icmp ult i8 [[TMP6]], 10
+// 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: [[TMP7:%.*]] = add i8 [[TMP5]], -97
-// DEFAULT-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP7]], 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: [[TMP8:%.*]] = add i8 [[TMP5]], -65
-// DEFAULT-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP8]], 6
+// 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 [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -87, %[[IF_ELSE_I_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 [[TMP5]] to i64
+// 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_I34_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I28_I_I10]], i64 1
-// DEFAULT-NEXT: [[TMP9]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// DEFAULT-NEXT: [[CMP_NOT_I30_I_I:%.*]] = icmp eq i8 [[TMP9]], 0
+// DEFAULT-NEXT: [[TMP7]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP13]]
// DEFAULT: [[WHILE_BODY_I_I_I]]:
-// DEFAULT-NEXT: [[TMP10:%.*]] = phi i8 [ [[TMP12:%.*]], %[[IF_THEN_I_I_I:.*]] ], [ [[TMP3]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// 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: [[TMP11:%.*]] = and i8 [[TMP10]], -8
-// DEFAULT-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP11]], 48
+// 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_I16]], 3
-// DEFAULT-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP10]] to i64
+// 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_I15]], i64 1
-// DEFAULT-NEXT: [[TMP12]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// DEFAULT-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP12]], 0
+// DEFAULT-NEXT: [[TMP10]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP9]]
// DEFAULT: [[WHILE_BODY_I18_I_I]]:
-// DEFAULT-NEXT: [[TMP13:%.*]] = phi i8 [ [[TMP15:%.*]], %[[IF_THEN_I21_I_I:.*]] ], [ [[TMP1]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// 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: [[TMP14:%.*]] = add i8 [[TMP13]], -48
-// DEFAULT-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP14]], 10
+// 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 [[TMP13]] to i64
+// 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: [[TMP15]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// DEFAULT-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP15]], 0
+// DEFAULT-NEXT: [[TMP13]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP12]]
// DEFAULT: [[_ZL3NANPKC_EXIT]]:
// DEFAULT-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_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: [[TMP16:%.*]] = bitcast i64 [[BF_SET9_I]] to double
-// DEFAULT-NEXT: ret double [[TMP16]]
+// 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]] {
@@ -4722,81 +4707,78 @@ extern "C" __device__ float test_nanf(const char *tag) {
// APPROX: [[IF_THEN_I_I]]:
// APPROX-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[TAG]], i64 1
// APPROX-NEXT: [[TMP2:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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_I_I_I_PREHEADER]]:
-// APPROX-NEXT: [[TMP3:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// APPROX-NEXT: [[CMP_NOT_I_I_I14:%.*]] = icmp eq i8 [[TMP3]], 0
// 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: [[TMP4:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// APPROX-NEXT: [[CMP_NOT_I30_I_I9:%.*]] = icmp eq i8 [[TMP4]], 0
-// APPROX-NEXT: br i1 [[CMP_NOT_I30_I_I9]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I31_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: [[TMP5:%.*]] = phi i8 [ [[TMP9:%.*]], %[[IF_END31_I_I_I:.*]] ], [ [[TMP4]], %[[IF_THEN5_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: [[TMP6:%.*]] = add i8 [[TMP5]], -48
-// APPROX-NEXT: [[OR_COND_I32_I_I:%.*]] = icmp ult i8 [[TMP6]], 10
+// 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: [[TMP7:%.*]] = add i8 [[TMP5]], -97
-// APPROX-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP7]], 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: [[TMP8:%.*]] = add i8 [[TMP5]], -65
-// APPROX-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP8]], 6
+// 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 [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -87, %[[IF_ELSE_I_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 [[TMP5]] to i64
+// 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_I34_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I28_I_I10]], i64 1
-// APPROX-NEXT: [[TMP9]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// APPROX-NEXT: [[CMP_NOT_I30_I_I:%.*]] = icmp eq i8 [[TMP9]], 0
+// APPROX-NEXT: [[TMP7]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP13]]
// APPROX: [[WHILE_BODY_I_I_I]]:
-// APPROX-NEXT: [[TMP10:%.*]] = phi i8 [ [[TMP12:%.*]], %[[IF_THEN_I_I_I:.*]] ], [ [[TMP3]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// 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: [[TMP11:%.*]] = and i8 [[TMP10]], -8
-// APPROX-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP11]], 48
+// 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_I16]], 3
-// APPROX-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP10]] to i64
+// 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_I15]], i64 1
-// APPROX-NEXT: [[TMP12]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// APPROX-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP12]], 0
+// APPROX-NEXT: [[TMP10]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP9]]
// APPROX: [[WHILE_BODY_I18_I_I]]:
-// APPROX-NEXT: [[TMP13:%.*]] = phi i8 [ [[TMP15:%.*]], %[[IF_THEN_I21_I_I:.*]] ], [ [[TMP1]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// 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: [[TMP14:%.*]] = add i8 [[TMP13]], -48
-// APPROX-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP14]], 10
+// 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 [[TMP13]] to i64
+// 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: [[TMP15]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// APPROX-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP15]], 0
+// APPROX-NEXT: [[TMP13]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP12]]
// APPROX: [[_ZL3NANPKC_EXIT]]:
// APPROX-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_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: [[TMP16:%.*]] = bitcast i64 [[BF_SET9_I]] to double
-// APPROX-NEXT: ret double [[TMP16]]
+// 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]] {
@@ -4811,81 +4793,78 @@ extern "C" __device__ float test_nanf(const char *tag) {
// NCRDIV: [[IF_THEN_I_I]]:
// NCRDIV-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds nuw i8, ptr [[TAG]], i64 1
// NCRDIV-NEXT: [[TMP2:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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_I_I_I_PREHEADER]]:
-// NCRDIV-NEXT: [[TMP3:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// NCRDIV-NEXT: [[CMP_NOT_I_I_I14:%.*]] = icmp eq i8 [[TMP3]], 0
// 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: [[TMP4:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// NCRDIV-NEXT: [[CMP_NOT_I30_I_I9:%.*]] = icmp eq i8 [[TMP4]], 0
-// NCRDIV-NEXT: br i1 [[CMP_NOT_I30_I_I9]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I31_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: [[TMP5:%.*]] = phi i8 [ [[TMP9:%.*]], %[[IF_END31_I_I_I:.*]] ], [ [[TMP4]], %[[IF_THEN5_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: [[TMP6:%.*]] = add i8 [[TMP5]], -48
-// NCRDIV-NEXT: [[OR_COND_I32_I_I:%.*]] = icmp ult i8 [[TMP6]], 10
+// 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: [[TMP7:%.*]] = add i8 [[TMP5]], -97
-// NCRDIV-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP7]], 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: [[TMP8:%.*]] = add i8 [[TMP5]], -65
-// NCRDIV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP8]], 6
+// 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 [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -87, %[[IF_ELSE_I_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 [[TMP5]] to i64
+// 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_I34_I_I]] = getelementptr inbounds nuw i8, ptr [[__TAGP_ADDR_0_I28_I_I10]], i64 1
-// NCRDIV-NEXT: [[TMP9]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// NCRDIV-NEXT: [[CMP_NOT_I30_I_I:%.*]] = icmp eq i8 [[TMP9]], 0
+// NCRDIV-NEXT: [[TMP7]] = load i8, ptr [[INCDEC_PTR_I34_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP13]]
// NCRDIV: [[WHILE_BODY_I_I_I]]:
-// NCRDIV-NEXT: [[TMP10:%.*]] = phi i8 [ [[TMP12:%.*]], %[[IF_THEN_I_I_I:.*]] ], [ [[TMP3]], %[[WHILE_COND_I_I_I_PREHEADER]] ]
+// 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: [[TMP11:%.*]] = and i8 [[TMP10]], -8
-// NCRDIV-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP11]], 48
+// 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_I16]], 3
-// NCRDIV-NEXT: [[CONV5_I_I_I:%.*]] = zext nneg i8 [[TMP10]] to i64
+// 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_I15]], i64 1
-// NCRDIV-NEXT: [[TMP12]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// NCRDIV-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP12]], 0
+// NCRDIV-NEXT: [[TMP10]] = load i8, ptr [[INCDEC_PTR_I_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP9]]
// NCRDIV: [[WHILE_BODY_I18_I_I]]:
-// NCRDIV-NEXT: [[TMP13:%.*]] = phi i8 [ [[TMP15:%.*]], %[[IF_THEN_I21_I_I:.*]] ], [ [[TMP1]], %[[WHILE_COND_I14_I_I_PREHEADER]] ]
+// 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: [[TMP14:%.*]] = add i8 [[TMP13]], -48
-// NCRDIV-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP14]], 10
+// 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 [[TMP13]] to i64
+// 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: [[TMP15]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
-// NCRDIV-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP15]], 0
+// NCRDIV-NEXT: [[TMP13]] = load i8, ptr [[INCDEC_PTR_I26_I_I]], align 1, !tbaa [[CHAR_TBAA8]]
+// 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 [[LOOP12]]
// NCRDIV: [[_ZL3NANPKC_EXIT]]:
// NCRDIV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_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: [[TMP16:%.*]] = bitcast i64 [[BF_SET9_I]] to double
-// NCRDIV-NEXT: ret double [[TMP16]]
+// 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]] {
More information about the cfe-commits
mailing list