// CHECK-NEXT: entry:
// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[P:%.*]], align 1, !tbaa [[TBAA3]]
// CHECK-NEXT: [[CMP_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// CHECK-NEXT: br i1 [[CMP_I]], label [[IF_THEN_I:%.*]], label [[WHILE_COND_I33_I:%.*]]
+// CHECK-NEXT: br i1 [[CMP_I]], label [[IF_THEN_I:%.*]], label [[WHILE_COND_I17_I:%.*]]
// CHECK: if.then.i:
// CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds i8, ptr [[P]], i64 1
// CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I17_I:%.*]] [
-// CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_PREHEADER:%.*]]
-// CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_PREHEADER]]
+// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I:%.*]] [
+// CHECK-NEXT: i8 120, label [[WHILE_COND_I33_I_PREHEADER:%.*]]
+// CHECK-NEXT: i8 88, label [[WHILE_COND_I33_I_PREHEADER]]
// CHECK-NEXT: ]
-// CHECK: while.cond.i.i.preheader:
-// CHECK-NEXT: br label [[WHILE_COND_I_I:%.*]]
-// CHECK: while.cond.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I:%.*]], [[CLEANUP_I_I:%.*]] ], [ [[INCDEC_PTR_I]], [[WHILE_COND_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[__R_0_I_I:%.*]] = phi i64 [ [[__R_2_I_I:%.*]], [[CLEANUP_I_I]] ], [ 0, [[WHILE_COND_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: [[CMP_NOT_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT:%.*]], label [[WHILE_BODY_I_I:%.*]]
-// CHECK: while.body.i.i:
+// CHECK: while.cond.i33.i.preheader:
+// CHECK-NEXT: br label [[WHILE_COND_I33_I:%.*]]
+// CHECK: while.cond.i33.i:
+// CHECK-NEXT: [[__TAGP_ADDR_0_I30_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I43_I:%.*]], [[CLEANUP_I44_I:%.*]] ], [ [[INCDEC_PTR_I]], [[WHILE_COND_I33_I_PREHEADER]] ]
+// CHECK-NEXT: [[__R_0_I31_I:%.*]] = phi i64 [ [[__R_2_I_I:%.*]], [[CLEANUP_I44_I]] ], [ 0, [[WHILE_COND_I33_I_PREHEADER]] ]
+// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT: [[CMP_NOT_I32_I:%.*]] = icmp eq i8 [[TMP2]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I32_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT:%.*]], label [[WHILE_BODY_I35_I:%.*]]
+// CHECK: while.body.i35.i:
// CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// CHECK-NEXT: [[OR_COND_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I_I]], label [[IF_END31_I_I:%.*]], label [[IF_ELSE_I_I:%.*]]
+// CHECK-NEXT: [[OR_COND_I34_I:%.*]] = icmp ult i8 [[TMP3]], 10
+// CHECK-NEXT: br i1 [[OR_COND_I34_I]], label [[IF_END31_I_I:%.*]], label [[IF_ELSE_I_I:%.*]]
// CHECK: if.else.i.i:
// CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
// CHECK-NEXT: [[OR_COND33_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
// CHECK: if.else17.i.i:
// CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
// CHECK-NEXT: [[OR_COND34_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// CHECK-NEXT: br i1 [[OR_COND34_I_I]], label [[IF_END31_I_I]], label [[CLEANUP_I_I]]
+// CHECK-NEXT: br i1 [[OR_COND34_I_I]], label [[IF_END31_I_I]], label [[CLEANUP_I44_I]]
// CHECK: if.end31.i.i:
-// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I]] ], [ -87, [[IF_ELSE_I_I]] ], [ -55, [[IF_ELSE17_I_I]] ]
-// CHECK-NEXT: [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I_I]], 4
+// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I35_I]] ], [ -87, [[IF_ELSE_I_I]] ], [ -55, [[IF_ELSE17_I_I]] ]
+// CHECK-NEXT: [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I31_I]], 4
// CHECK-NEXT: [[CONV25_I_I:%.*]] = sext i8 [[TMP2]] to i64
// CHECK-NEXT: [[ADD26_I_I:%.*]] = add i64 [[MUL24_I_I]], [[DOTSINK]]
// CHECK-NEXT: [[ADD28_I_I:%.*]] = add i64 [[ADD26_I_I]], [[CONV25_I_I]]
+// CHECK-NEXT: [[INCDEC_PTR_I42_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I]], i64 1
+// CHECK-NEXT: br label [[CLEANUP_I44_I]]
+// CHECK: cleanup.i44.i:
+// CHECK-NEXT: [[__TAGP_ADDR_1_I43_I]] = phi ptr [ [[INCDEC_PTR_I42_I]], [[IF_END31_I_I]] ], [ [[__TAGP_ADDR_0_I30_I]], [[IF_ELSE17_I_I]] ]
+// CHECK-NEXT: [[__R_2_I_I]] = phi i64 [ [[ADD28_I_I]], [[IF_END31_I_I]] ], [ [[__R_0_I31_I]], [[IF_ELSE17_I_I]] ]
+// CHECK-NEXT: [[COND_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I]] ], [ false, [[IF_ELSE17_I_I]] ]
+// CHECK-NEXT: br i1 [[COND_I_I]], label [[WHILE_COND_I33_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP10]]
+// CHECK: while.cond.i.i:
+// CHECK-NEXT: [[__TAGP_ADDR_0_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I:%.*]], [[CLEANUP_I_I:%.*]] ], [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ]
+// CHECK-NEXT: [[__R_0_I_I:%.*]] = phi i64 [ [[__R_1_I_I:%.*]], [[CLEANUP_I_I]] ], [ 0, [[IF_THEN_I]] ]
+// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT: [[CMP_NOT_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I_I:%.*]]
+// CHECK: while.body.i.i:
+// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// CHECK-NEXT: [[OR_COND_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// CHECK-NEXT: br i1 [[OR_COND_I_I]], label [[IF_THEN_I_I:%.*]], label [[CLEANUP_I_I]]
+// CHECK: if.then.i.i:
+// CHECK-NEXT: [[MUL_I_I:%.*]] = shl i64 [[__R_0_I_I]], 3
+// CHECK-NEXT: [[CONV5_I_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT: [[ADD_I_I:%.*]] = add i64 [[MUL_I_I]], -48
+// CHECK-NEXT: [[SUB_I_I:%.*]] = add i64 [[ADD_I_I]], [[CONV5_I_I]]
// CHECK-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I]], i64 1
// CHECK-NEXT: br label [[CLEANUP_I_I]]
// CHECK: cleanup.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I_I]] = phi ptr [ [[INCDEC_PTR_I_I]], [[IF_END31_I_I]] ], [ [[__TAGP_ADDR_0_I_I]], [[IF_ELSE17_I_I]] ]
-// CHECK-NEXT: [[__R_2_I_I]] = phi i64 [ [[ADD28_I_I]], [[IF_END31_I_I]] ], [ [[__R_0_I_I]], [[IF_ELSE17_I_I]] ]
-// CHECK-NEXT: [[COND_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I]] ], [ false, [[IF_ELSE17_I_I]] ]
-// CHECK-NEXT: br i1 [[COND_I_I]], label [[WHILE_COND_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP10]]
+// CHECK-NEXT: [[__TAGP_ADDR_1_I_I]] = phi ptr [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ], [ [[__TAGP_ADDR_0_I_I]], [[WHILE_BODY_I_I]] ]
+// CHECK-NEXT: [[__R_1_I_I]] = phi i64 [ [[SUB_I_I]], [[IF_THEN_I_I]] ], [ [[__R_0_I_I]], [[WHILE_BODY_I_I]] ]
+// CHECK-NEXT: br i1 [[OR_COND_I_I]], label [[WHILE_COND_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP6]]
// CHECK: while.cond.i17.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I14_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I:%.*]], [[CLEANUP_I28_I:%.*]] ], [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ]
-// CHECK-NEXT: [[__R_0_I15_I:%.*]] = phi i64 [ [[__R_1_I27_I:%.*]], [[CLEANUP_I28_I]] ], [ 0, [[IF_THEN_I]] ]
-// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: [[CMP_NOT_I16_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// CHECK-NEXT: [[__TAGP_ADDR_0_I14_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I:%.*]], [[CLEANUP_I28_I:%.*]] ], [ [[P]], [[ENTRY:%.*]] ]
+// CHECK-NEXT: [[__R_0_I15_I:%.*]] = phi i64 [ [[__R_1_I27_I:%.*]], [[CLEANUP_I28_I]] ], [ 0, [[ENTRY]] ]
+// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT: [[CMP_NOT_I16_I:%.*]] = icmp eq i8 [[TMP8]], 0
// CHECK-NEXT: br i1 [[CMP_NOT_I16_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I19_I:%.*]]
// CHECK: while.body.i19.i:
-// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// CHECK-NEXT: [[OR_COND_I18_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// CHECK-NEXT: [[OR_COND_I18_I:%.*]] = icmp ult i8 [[TMP9]], 10
// CHECK-NEXT: br i1 [[OR_COND_I18_I]], label [[IF_THEN_I25_I:%.*]], label [[CLEANUP_I28_I]]
// CHECK: if.then.i25.i:
-// CHECK-NEXT: [[MUL_I20_I:%.*]] = shl i64 [[__R_0_I15_I]], 3
-// CHECK-NEXT: [[CONV5_I21_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT: [[MUL_I20_I:%.*]] = mul i64 [[__R_0_I15_I]], 10
+// CHECK-NEXT: [[CONV5_I21_I:%.*]] = sext i8 [[TMP8]] to i64
// CHECK-NEXT: [[ADD_I22_I:%.*]] = add i64 [[MUL_I20_I]], -48
// CHECK-NEXT: [[SUB_I23_I:%.*]] = add i64 [[ADD_I22_I]], [[CONV5_I21_I]]
// CHECK-NEXT: [[INCDEC_PTR_I24_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I14_I]], i64 1
// CHECK: cleanup.i28.i:
// CHECK-NEXT: [[__TAGP_ADDR_1_I26_I]] = phi ptr [ [[INCDEC_PTR_I24_I]], [[IF_THEN_I25_I]] ], [ [[__TAGP_ADDR_0_I14_I]], [[WHILE_BODY_I19_I]] ]
// CHECK-NEXT: [[__R_1_I27_I]] = phi i64 [ [[SUB_I23_I]], [[IF_THEN_I25_I]] ], [ [[__R_0_I15_I]], [[WHILE_BODY_I19_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I18_I]], label [[WHILE_COND_I17_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP6]]
-// CHECK: while.cond.i33.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I30_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I42_I:%.*]], [[CLEANUP_I44_I:%.*]] ], [ [[P]], [[ENTRY:%.*]] ]
-// CHECK-NEXT: [[__R_0_I31_I:%.*]] = phi i64 [ [[__R_1_I43_I:%.*]], [[CLEANUP_I44_I]] ], [ 0, [[ENTRY]] ]
-// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: [[CMP_NOT_I32_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I32_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I35_I:%.*]]
-// CHECK: while.body.i35.i:
-// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// CHECK-NEXT: [[OR_COND_I34_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I34_I]], label [[IF_THEN_I41_I:%.*]], label [[CLEANUP_I44_I]]
-// CHECK: if.then.i41.i:
-// CHECK-NEXT: [[MUL_I36_I:%.*]] = mul i64 [[__R_0_I31_I]], 10
-// CHECK-NEXT: [[CONV5_I37_I:%.*]] = sext i8 [[TMP8]] to i64
-// CHECK-NEXT: [[ADD_I38_I:%.*]] = add i64 [[MUL_I36_I]], -48
-// CHECK-NEXT: [[SUB_I39_I:%.*]] = add i64 [[ADD_I38_I]], [[CONV5_I37_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I40_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I]], i64 1
-// CHECK-NEXT: br label [[CLEANUP_I44_I]]
-// CHECK: cleanup.i44.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I42_I]] = phi ptr [ [[INCDEC_PTR_I40_I]], [[IF_THEN_I41_I]] ], [ [[__TAGP_ADDR_0_I30_I]], [[WHILE_BODY_I35_I]] ]
-// CHECK-NEXT: [[__R_1_I43_I]] = phi i64 [ [[SUB_I39_I]], [[IF_THEN_I41_I]] ], [ [[__R_0_I31_I]], [[WHILE_BODY_I35_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I34_I]], label [[WHILE_COND_I33_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP9]]
+// CHECK-NEXT: br i1 [[OR_COND_I18_I]], label [[WHILE_COND_I17_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP9]]
// CHECK: _ZL15__make_mantissaPKc.exit:
-// CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, [[CLEANUP_I28_I]] ], [ [[__R_0_I15_I]], [[WHILE_COND_I17_I]] ], [ 0, [[CLEANUP_I_I]] ], [ [[__R_0_I_I]], [[WHILE_COND_I_I]] ], [ 0, [[CLEANUP_I44_I]] ], [ [[__R_0_I31_I]], [[WHILE_COND_I33_I]] ]
+// CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I]] ], [ [[__R_0_I_I]], [[WHILE_COND_I_I]] ], [ 0, [[CLEANUP_I44_I]] ], [ [[__R_0_I31_I]], [[WHILE_COND_I33_I]] ], [ 0, [[CLEANUP_I28_I]] ], [ [[__R_0_I15_I]], [[WHILE_COND_I17_I]] ]
// CHECK-NEXT: ret i64 [[RETVAL_0_I]]
//
extern "C" __device__ uint64_t test___make_mantissa(const char *p) {
return exp2f(x);
}
-//
// DEFAULT-LABEL: @test_exp2(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// DEFAULT-NEXT: ]
// DEFAULT: if.then.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL3JNFIF_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
// DEFAULT: for.body.i:
-// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
-// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]]
-// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]]
-// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP13:![0-9]+]]
+// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
+// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
+// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// DEFAULT-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP13:![0-9]+]]
// DEFAULT: _ZL3jnfif.exit:
-// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// DEFAULT-NEXT: ret float [[RETVAL_0_I]]
//
// FINITEONLY-LABEL: @test_jnf(
// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// FINITEONLY-NEXT: ]
// FINITEONLY: if.then.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
// FINITEONLY: for.body.i:
-// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]]
-// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_025_I]], [[DIV_I]]
-// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_024_I]]
-// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP13:![0-9]+]]
+// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_0_I3]], [[DIV_I]]
+// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_0_I2]]
+// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// FINITEONLY-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP13:![0-9]+]]
// FINITEONLY: _ZL3jnfif.exit:
-// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// FINITEONLY-NEXT: ret float [[RETVAL_0_I]]
//
extern "C" __device__ float test_jnf(int x, float y) {
// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// DEFAULT-NEXT: ]
// DEFAULT: if.then.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL2JNID_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL2JNID_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
// DEFAULT: for.body.i:
-// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
-// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]]
-// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]]
-// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
+// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
+// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]
+// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// DEFAULT-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
// DEFAULT: _ZL2jnid.exit:
-// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// DEFAULT-NEXT: ret double [[RETVAL_0_I]]
//
// FINITEONLY-LABEL: @test_jn(
// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// FINITEONLY-NEXT: ]
// FINITEONLY: if.then.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
// FINITEONLY: for.body.i:
-// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]]
-// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_025_I]], [[DIV_I]]
-// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_024_I]]
-// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
+// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_0_I3]], [[DIV_I]]
+// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_0_I2]]
+// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// FINITEONLY-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
// FINITEONLY: _ZL2jnid.exit:
-// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// FINITEONLY-NEXT: ret double [[RETVAL_0_I]]
//
extern "C" __device__ double test_jn(int x, double y) {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa [[TBAA3]]
// CHECK-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I33_I_I:%.*]]
+// CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I17_I_I:%.*]]
// CHECK: if.then.i.i:
// CHECK-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[TAG]], i64 1
// CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I17_I_I:%.*]] [
-// CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_I_PREHEADER:%.*]]
-// CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_I_PREHEADER]]
+// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [
+// CHECK-NEXT: i8 120, label [[WHILE_COND_I33_I_I_PREHEADER:%.*]]
+// CHECK-NEXT: i8 88, label [[WHILE_COND_I33_I_I_PREHEADER]]
// CHECK-NEXT: ]
-// CHECK: while.cond.i.i.i.preheader:
-// CHECK-NEXT: br label [[WHILE_COND_I_I_I:%.*]]
-// CHECK: while.cond.i.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[WHILE_COND_I_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL4NANFPKC_EXIT:%.*]], label [[WHILE_BODY_I_I_I:%.*]]
-// CHECK: while.body.i.i.i:
+// CHECK: while.cond.i33.i.i.preheader:
+// CHECK-NEXT: br label [[WHILE_COND_I33_I_I:%.*]]
+// CHECK: while.cond.i33.i.i:
+// CHECK-NEXT: [[__TAGP_ADDR_0_I30_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I43_I_I:%.*]], [[CLEANUP_I44_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I33_I_I_PREHEADER]] ]
+// CHECK-NEXT: [[__R_0_I31_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I44_I_I]] ], [ 0, [[WHILE_COND_I33_I_I_PREHEADER]] ]
+// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT: [[CMP_NOT_I32_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I32_I_I]], label [[_ZL4NANFPKC_EXIT:%.*]], label [[WHILE_BODY_I35_I_I:%.*]]
+// CHECK: while.body.i35.i.i:
// CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// CHECK-NEXT: [[OR_COND_I_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
+// CHECK-NEXT: [[OR_COND_I34_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
+// CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
// CHECK: if.else.i.i.i:
// CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
// CHECK-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
// CHECK: if.else17.i.i.i:
// CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
// CHECK-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// CHECK-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I_I_I]]
+// CHECK-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I44_I_I]]
// CHECK: if.end31.i.i.i:
-// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 4
+// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I35_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I31_I_I]], 4
// CHECK-NEXT: [[CONV25_I_I_I:%.*]] = sext i8 [[TMP2]] to i64
// CHECK-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
// CHECK-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
+// CHECK-NEXT: [[INCDEC_PTR_I42_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I_I]], i64 1
+// CHECK-NEXT: br label [[CLEANUP_I44_I_I]]
+// CHECK: cleanup.i44.i.i:
+// CHECK-NEXT: [[__TAGP_ADDR_1_I43_I_I]] = phi ptr [ [[INCDEC_PTR_I42_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I30_I_I]], [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I31_I_I]], [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I33_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP10]]
+// CHECK: while.cond.i.i.i:
+// CHECK-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
+// CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
+// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I_I_I:%.*]]
+// CHECK: while.body.i.i.i:
+// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// CHECK-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_THEN_I_I_I:%.*]], label [[CLEANUP_I_I_I]]
+// CHECK: if.then.i.i.i:
+// CHECK-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
+// CHECK-NEXT: [[CONV5_I_I_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
+// CHECK-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
// CHECK-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1
// CHECK-NEXT: br label [[CLEANUP_I_I_I]]
// CHECK: cleanup.i.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I_I_I]], [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP10]]
+// CHECK-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// CHECK-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP6]]
// CHECK: while.cond.i17.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I_I:%.*]], [[CLEANUP_I28_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
-// CHECK-NEXT: [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I27_I_I:%.*]], [[CLEANUP_I28_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
-// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// CHECK-NEXT: [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I_I:%.*]], [[CLEANUP_I28_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
+// CHECK-NEXT: [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I27_I_I:%.*]], [[CLEANUP_I28_I_I]] ], [ 0, [[ENTRY]] ]
+// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT: [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
// CHECK-NEXT: br i1 [[CMP_NOT_I16_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I19_I_I:%.*]]
// CHECK: while.body.i19.i.i:
-// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// CHECK-NEXT: [[OR_COND_I18_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// CHECK-NEXT: [[OR_COND_I18_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
// CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[IF_THEN_I25_I_I:%.*]], label [[CLEANUP_I28_I_I]]
// CHECK: if.then.i25.i.i:
-// CHECK-NEXT: [[MUL_I20_I_I:%.*]] = shl i64 [[__R_0_I15_I_I]], 3
-// CHECK-NEXT: [[CONV5_I21_I_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT: [[MUL_I20_I_I:%.*]] = mul i64 [[__R_0_I15_I_I]], 10
+// CHECK-NEXT: [[CONV5_I21_I_I:%.*]] = sext i8 [[TMP8]] to i64
// CHECK-NEXT: [[ADD_I22_I_I:%.*]] = add i64 [[MUL_I20_I_I]], -48
// CHECK-NEXT: [[SUB_I23_I_I:%.*]] = add i64 [[ADD_I22_I_I]], [[CONV5_I21_I_I]]
// CHECK-NEXT: [[INCDEC_PTR_I24_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I14_I_I]], i64 1
// CHECK: cleanup.i28.i.i:
// CHECK-NEXT: [[__TAGP_ADDR_1_I26_I_I]] = phi ptr [ [[INCDEC_PTR_I24_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__TAGP_ADDR_0_I14_I_I]], [[WHILE_BODY_I19_I_I]] ]
// CHECK-NEXT: [[__R_1_I27_I_I]] = phi i64 [ [[SUB_I23_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_BODY_I19_I_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[WHILE_COND_I17_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP6]]
-// CHECK: while.cond.i33.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I30_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I42_I_I:%.*]], [[CLEANUP_I44_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
-// CHECK-NEXT: [[__R_0_I31_I_I:%.*]] = phi i64 [ [[__R_1_I43_I_I:%.*]], [[CLEANUP_I44_I_I]] ], [ 0, [[ENTRY]] ]
-// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: [[CMP_NOT_I32_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I32_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I35_I_I:%.*]]
-// CHECK: while.body.i35.i.i:
-// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// CHECK-NEXT: [[OR_COND_I34_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[IF_THEN_I41_I_I:%.*]], label [[CLEANUP_I44_I_I]]
-// CHECK: if.then.i41.i.i:
-// CHECK-NEXT: [[MUL_I36_I_I:%.*]] = mul i64 [[__R_0_I31_I_I]], 10
-// CHECK-NEXT: [[CONV5_I37_I_I:%.*]] = sext i8 [[TMP8]] to i64
-// CHECK-NEXT: [[ADD_I38_I_I:%.*]] = add i64 [[MUL_I36_I_I]], -48
-// CHECK-NEXT: [[SUB_I39_I_I:%.*]] = add i64 [[ADD_I38_I_I]], [[CONV5_I37_I_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I_I]], i64 1
-// CHECK-NEXT: br label [[CLEANUP_I44_I_I]]
-// CHECK: cleanup.i44.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I42_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__TAGP_ADDR_0_I30_I_I]], [[WHILE_BODY_I35_I_I]] ]
-// CHECK-NEXT: [[__R_1_I43_I_I]] = phi i64 [ [[SUB_I39_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_BODY_I35_I_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[WHILE_COND_I33_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP9]]
+// CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[WHILE_COND_I17_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP9]]
// CHECK: _ZL4nanfPKc.exit:
-// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I28_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ], [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I44_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_COND_I33_I_I]] ]
+// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I44_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_COND_I33_I_I]] ], [ 0, [[CLEANUP_I28_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ]
// CHECK-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32
// CHECK-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303
// CHECK-NEXT: [[BF_SET9_I:%.*]] = or i32 [[BF_VALUE_I]], 2143289344
// CHECK-NEXT: entry:
// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa [[TBAA3]]
// CHECK-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
-// CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I33_I_I:%.*]]
+// CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I17_I_I:%.*]]
// CHECK: if.then.i.i:
// CHECK-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[TAG]], i64 1
// CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I17_I_I:%.*]] [
-// CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_I_PREHEADER:%.*]]
-// CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_I_PREHEADER]]
+// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [
+// CHECK-NEXT: i8 120, label [[WHILE_COND_I33_I_I_PREHEADER:%.*]]
+// CHECK-NEXT: i8 88, label [[WHILE_COND_I33_I_I_PREHEADER]]
// CHECK-NEXT: ]
-// CHECK: while.cond.i.i.i.preheader:
-// CHECK-NEXT: br label [[WHILE_COND_I_I_I:%.*]]
-// CHECK: while.cond.i.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[WHILE_COND_I_I_I_PREHEADER]] ]
-// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL3NANPKC_EXIT:%.*]], label [[WHILE_BODY_I_I_I:%.*]]
-// CHECK: while.body.i.i.i:
+// CHECK: while.cond.i33.i.i.preheader:
+// CHECK-NEXT: br label [[WHILE_COND_I33_I_I:%.*]]
+// CHECK: while.cond.i33.i.i:
+// CHECK-NEXT: [[__TAGP_ADDR_0_I30_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I43_I_I:%.*]], [[CLEANUP_I44_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I33_I_I_PREHEADER]] ]
+// CHECK-NEXT: [[__R_0_I31_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I44_I_I]] ], [ 0, [[WHILE_COND_I33_I_I_PREHEADER]] ]
+// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT: [[CMP_NOT_I32_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I32_I_I]], label [[_ZL3NANPKC_EXIT:%.*]], label [[WHILE_BODY_I35_I_I:%.*]]
+// CHECK: while.body.i35.i.i:
// CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
-// CHECK-NEXT: [[OR_COND_I_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
+// CHECK-NEXT: [[OR_COND_I34_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
+// CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
// CHECK: if.else.i.i.i:
// CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
// CHECK-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
// CHECK: if.else17.i.i.i:
// CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
// CHECK-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
-// CHECK-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I_I_I]]
+// CHECK-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I44_I_I]]
// CHECK: if.end31.i.i.i:
-// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 4
+// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I35_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I31_I_I]], 4
// CHECK-NEXT: [[CONV25_I_I_I:%.*]] = sext i8 [[TMP2]] to i64
// CHECK-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
// CHECK-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
+// CHECK-NEXT: [[INCDEC_PTR_I42_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I_I]], i64 1
+// CHECK-NEXT: br label [[CLEANUP_I44_I_I]]
+// CHECK: cleanup.i44.i.i:
+// CHECK-NEXT: [[__TAGP_ADDR_1_I43_I_I]] = phi ptr [ [[INCDEC_PTR_I42_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I30_I_I]], [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I31_I_I]], [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ]
+// CHECK-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I33_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP10]]
+// CHECK: while.cond.i.i.i:
+// CHECK-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
+// CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
+// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// CHECK-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I_I_I:%.*]]
+// CHECK: while.body.i.i.i:
+// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
+// CHECK-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_THEN_I_I_I:%.*]], label [[CLEANUP_I_I_I]]
+// CHECK: if.then.i.i.i:
+// CHECK-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3
+// CHECK-NEXT: [[CONV5_I_I_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48
+// CHECK-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]]
// CHECK-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1
// CHECK-NEXT: br label [[CLEANUP_I_I_I]]
// CHECK: cleanup.i.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I_I_I]], [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ]
-// CHECK-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP10]]
+// CHECK-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// CHECK-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_BODY_I_I_I]] ]
+// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP6]]
// CHECK: while.cond.i17.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I_I:%.*]], [[CLEANUP_I28_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
-// CHECK-NEXT: [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I27_I_I:%.*]], [[CLEANUP_I28_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
-// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
+// CHECK-NEXT: [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I_I:%.*]], [[CLEANUP_I28_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
+// CHECK-NEXT: [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I27_I_I:%.*]], [[CLEANUP_I28_I_I]] ], [ 0, [[ENTRY]] ]
+// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT: [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
// CHECK-NEXT: br i1 [[CMP_NOT_I16_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I19_I_I:%.*]]
// CHECK: while.body.i19.i.i:
-// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
-// CHECK-NEXT: [[OR_COND_I18_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
+// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
+// CHECK-NEXT: [[OR_COND_I18_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
// CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[IF_THEN_I25_I_I:%.*]], label [[CLEANUP_I28_I_I]]
// CHECK: if.then.i25.i.i:
-// CHECK-NEXT: [[MUL_I20_I_I:%.*]] = shl i64 [[__R_0_I15_I_I]], 3
-// CHECK-NEXT: [[CONV5_I21_I_I:%.*]] = sext i8 [[TMP6]] to i64
+// CHECK-NEXT: [[MUL_I20_I_I:%.*]] = mul i64 [[__R_0_I15_I_I]], 10
+// CHECK-NEXT: [[CONV5_I21_I_I:%.*]] = sext i8 [[TMP8]] to i64
// CHECK-NEXT: [[ADD_I22_I_I:%.*]] = add i64 [[MUL_I20_I_I]], -48
// CHECK-NEXT: [[SUB_I23_I_I:%.*]] = add i64 [[ADD_I22_I_I]], [[CONV5_I21_I_I]]
// CHECK-NEXT: [[INCDEC_PTR_I24_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I14_I_I]], i64 1
// CHECK: cleanup.i28.i.i:
// CHECK-NEXT: [[__TAGP_ADDR_1_I26_I_I]] = phi ptr [ [[INCDEC_PTR_I24_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__TAGP_ADDR_0_I14_I_I]], [[WHILE_BODY_I19_I_I]] ]
// CHECK-NEXT: [[__R_1_I27_I_I]] = phi i64 [ [[SUB_I23_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_BODY_I19_I_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[WHILE_COND_I17_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP6]]
-// CHECK: while.cond.i33.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_0_I30_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I42_I_I:%.*]], [[CLEANUP_I44_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
-// CHECK-NEXT: [[__R_0_I31_I_I:%.*]] = phi i64 [ [[__R_1_I43_I_I:%.*]], [[CLEANUP_I44_I_I]] ], [ 0, [[ENTRY]] ]
-// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I_I]], align 1, !tbaa [[TBAA3]]
-// CHECK-NEXT: [[CMP_NOT_I32_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
-// CHECK-NEXT: br i1 [[CMP_NOT_I32_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I35_I_I:%.*]]
-// CHECK: while.body.i35.i.i:
-// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
-// CHECK-NEXT: [[OR_COND_I34_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
-// CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[IF_THEN_I41_I_I:%.*]], label [[CLEANUP_I44_I_I]]
-// CHECK: if.then.i41.i.i:
-// CHECK-NEXT: [[MUL_I36_I_I:%.*]] = mul i64 [[__R_0_I31_I_I]], 10
-// CHECK-NEXT: [[CONV5_I37_I_I:%.*]] = sext i8 [[TMP8]] to i64
-// CHECK-NEXT: [[ADD_I38_I_I:%.*]] = add i64 [[MUL_I36_I_I]], -48
-// CHECK-NEXT: [[SUB_I39_I_I:%.*]] = add i64 [[ADD_I38_I_I]], [[CONV5_I37_I_I]]
-// CHECK-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I_I]], i64 1
-// CHECK-NEXT: br label [[CLEANUP_I44_I_I]]
-// CHECK: cleanup.i44.i.i:
-// CHECK-NEXT: [[__TAGP_ADDR_1_I42_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__TAGP_ADDR_0_I30_I_I]], [[WHILE_BODY_I35_I_I]] ]
-// CHECK-NEXT: [[__R_1_I43_I_I]] = phi i64 [ [[SUB_I39_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_BODY_I35_I_I]] ]
-// CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[WHILE_COND_I33_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP9]]
+// CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[WHILE_COND_I17_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP9]]
// CHECK: _ZL3nanPKc.exit:
-// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I28_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ], [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I44_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_COND_I33_I_I]] ]
+// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I44_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_COND_I33_I_I]] ], [ 0, [[CLEANUP_I28_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ]
// CHECK-NEXT: [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]], 2251799813685247
// CHECK-NEXT: [[BF_SET9_I:%.*]] = or i64 [[BF_VALUE_I]], 9221120237041090560
// CHECK-NEXT: [[TMP10:%.*]] = bitcast i64 [[BF_SET9_I]] to double
// DEFAULT-LABEL: @test_normf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// DEFAULT: while.body.i:
-// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA15]]
+// DEFAULT-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA15]]
// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
-// DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]]
-// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1
// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP19:![0-9]+]]
// DEFAULT: _ZL5normfiPKf.exit:
-// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
+// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_normf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// FINITEONLY: while.body.i:
-// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA15]]
+// FINITEONLY-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA15]]
// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]]
-// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_08_I]], [[MUL_I]]
-// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_0_I4]], [[MUL_I]]
+// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1
// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP19:![0-9]+]]
// FINITEONLY: _ZL5normfiPKf.exit:
-// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
+// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_normf(int x, const float *y) {
// DEFAULT-LABEL: @test_norm(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// DEFAULT: while.body.i:
-// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA17]]
+// DEFAULT-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA17]]
// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
-// DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]]
-// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1
// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
// DEFAULT: _ZL4normiPKd.exit:
-// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
+// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_norm(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// FINITEONLY: while.body.i:
-// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA17]]
+// FINITEONLY-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA17]]
// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]]
-// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_08_I]], [[MUL_I]]
-// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_0_I4]], [[MUL_I]]
+// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1
// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
// FINITEONLY: _ZL4normiPKd.exit:
-// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
+// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR13]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_norm(int x, const double *y) {
// DEFAULT-LABEL: @test_rnormf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// DEFAULT: while.body.i:
-// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA15]]
+// DEFAULT-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA15]]
// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
-// DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]]
-// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1
// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
// DEFAULT: _ZL6rnormfiPKf.exit:
-// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR14]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rnormf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// FINITEONLY: while.body.i:
-// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA15]]
+// FINITEONLY-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA15]]
// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]]
-// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_08_I]], [[MUL_I]]
-// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_0_I4]], [[MUL_I]]
+// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1
// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
// FINITEONLY: _ZL6rnormfiPKf.exit:
-// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rsqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR14]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
//
extern "C" __device__ float test_rnormf(int x, const float* y) {
// DEFAULT-LABEL: @test_rnorm(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// DEFAULT: while.body.i:
-// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA17]]
+// DEFAULT-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA17]]
// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
-// DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]]
-// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
+// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1
// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
// DEFAULT: _ZL5rnormiPKd.exit:
-// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR14]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_rnorm(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
-// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
// FINITEONLY: while.body.i:
-// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
-// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
-// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
-// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA17]]
+// FINITEONLY-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1
+// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA17]]
// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]]
-// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_08_I]], [[MUL_I]]
-// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_0_I4]], [[MUL_I]]
+// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1
// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
// FINITEONLY: _ZL5rnormiPKd.exit:
-// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rsqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR14]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
//
extern "C" __device__ double test_rnorm(int x, const double* y) {
// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// DEFAULT-NEXT: ]
// DEFAULT: if.then.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL3YNFIF_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
// DEFAULT: for.body.i:
-// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
-// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]]
-// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]]
-// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
+// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
+// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
+// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// DEFAULT-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
// DEFAULT: _ZL3ynfif.exit:
-// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// DEFAULT-NEXT: ret float [[RETVAL_0_I]]
//
// FINITEONLY-LABEL: @test_ynf(
// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// FINITEONLY-NEXT: ]
// FINITEONLY: if.then.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
// FINITEONLY: for.body.i:
-// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]]
-// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_025_I]], [[DIV_I]]
-// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_024_I]]
-// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
+// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_0_I3]], [[DIV_I]]
+// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_0_I2]]
+// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// FINITEONLY-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
// FINITEONLY: _ZL3ynfif.exit:
-// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// FINITEONLY-NEXT: ret float [[RETVAL_0_I]]
//
extern "C" __device__ float test_ynf(int x, float y) {
// DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// DEFAULT-NEXT: ]
// DEFAULT: if.then.i:
-// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL2YNID_EXIT:%.*]]
// DEFAULT: if.then2.i:
-// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
// DEFAULT-NEXT: br label [[_ZL2YNID_EXIT]]
// DEFAULT: if.end4.i:
-// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
-// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
+// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
+// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
// DEFAULT: for.body.i:
-// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
-// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]]
-// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]]
-// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
+// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
+// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]
+// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// DEFAULT-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
// DEFAULT: _ZL2ynid.exit:
-// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// DEFAULT-NEXT: ret double [[RETVAL_0_I]]
//
// FINITEONLY-LABEL: @test_yn(
// FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
// FINITEONLY-NEXT: ]
// FINITEONLY: if.then.i:
-// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT:%.*]]
// FINITEONLY: if.then2.i:
-// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
// FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT]]
// FINITEONLY: if.end4.i:
-// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
-// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
-// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
+// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
+// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
// FINITEONLY: for.body.i:
-// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
-// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]]
-// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_025_I]], [[DIV_I]]
-// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_024_I]]
-// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
-// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
-// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
+// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_0_I3]], [[DIV_I]]
+// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_0_I2]]
+// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1
+// FINITEONLY-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
// FINITEONLY: _ZL2ynid.exit:
-// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
+// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
// FINITEONLY-NEXT: ret double [[RETVAL_0_I]]
//
extern "C" __device__ double test_yn(int x, double y) {