diff --git a/clang/test/Frontend/optimization-remark-with-hotness-new-pm.c b/clang/test/Frontend/optimization-remark-with-hotness-new-pm.c index cb1992ec4cb85..173b43ba41bfe 100644 --- a/clang/test/Frontend/optimization-remark-with-hotness-new-pm.c +++ b/clang/test/Frontend/optimization-remark-with-hotness-new-pm.c @@ -73,7 +73,7 @@ void bar(int x) { // THRESHOLD-NOT: hotness // NO_PGO: '-fdiagnostics-show-hotness' requires profile-guided optimization information // NO_PGO: '-fdiagnostics-hotness-threshold=' requires profile-guided optimization information - // expected-remark@+1 {{'foo' inlined into 'bar': always inline attribute at callsite bar:8:10; (hotness:}} + // expected-remark@+1 {{'foo' inlined into 'bar' with (cost=always): always inline attribute at callsite bar:8:10; (hotness:}} sum += foo(x, x - 2); } diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index fc18e14d82296..b57c38d70b14c 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -138,26 +138,26 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) { // CHECK-NEXT: entry: // CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[P:%.*]], align 1, !tbaa [[TBAA4]] // CHECK-NEXT: [[CMP_I:%.*]] = icmp eq i8 [[TMP0]], 48 -// CHECK-NEXT: br i1 [[CMP_I]], label [[IF_THEN_I:%.*]], label [[WHILE_COND_I30_I:%.*]] +// CHECK-NEXT: br i1 [[CMP_I]], label [[IF_THEN_I:%.*]], label [[WHILE_COND_I14_I:%.*]] // CHECK: if.then.i: // CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds i8, ptr [[P]], i64 1 // CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I]], align 1, !tbaa [[TBAA4]] -// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I14_I:%.*]] [ -// CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_PREHEADER:%.*]] -// CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_PREHEADER]] +// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I:%.*]] [ +// CHECK-NEXT: i8 120, label [[WHILE_COND_I30_I_PREHEADER:%.*]] +// CHECK-NEXT: i8 88, label [[WHILE_COND_I30_I_PREHEADER]] // CHECK-NEXT: ] -// CHECK: while.cond.i.i.preheader: -// CHECK-NEXT: br label [[WHILE_COND_I_I:%.*]] -// CHECK: while.cond.i.i: -// CHECK-NEXT: [[__TAGP_ADDR_0_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I:%.*]], [[CLEANUP_I_I:%.*]] ], [ [[INCDEC_PTR_I]], [[WHILE_COND_I_I_PREHEADER]] ] -// CHECK-NEXT: [[__R_0_I_I:%.*]] = phi i64 [ [[__R_2_I_I:%.*]], [[CLEANUP_I_I]] ], [ 0, [[WHILE_COND_I_I_PREHEADER]] ] -// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I]], align 1, !tbaa [[TBAA4]] -// CHECK-NEXT: [[CMP_NOT_I_I:%.*]] = icmp eq i8 [[TMP2]], 0 -// CHECK-NEXT: br i1 [[CMP_NOT_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT:%.*]], label [[WHILE_BODY_I_I:%.*]] -// CHECK: while.body.i.i: +// CHECK: while.cond.i30.i.preheader: +// CHECK-NEXT: br label [[WHILE_COND_I30_I:%.*]] +// CHECK: while.cond.i30.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I31_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I:%.*]], [[CLEANUP_I36_I:%.*]] ], [ [[INCDEC_PTR_I]], [[WHILE_COND_I30_I_PREHEADER]] ] +// CHECK-NEXT: [[__R_0_I32_I:%.*]] = phi i64 [ [[__R_2_I_I:%.*]], [[CLEANUP_I36_I]] ], [ 0, [[WHILE_COND_I30_I_PREHEADER]] ] +// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I]], align 1, !tbaa [[TBAA4]] +// CHECK-NEXT: [[CMP_NOT_I33_I:%.*]] = icmp eq i8 [[TMP2]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I33_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT:%.*]], label [[WHILE_BODY_I34_I:%.*]] +// CHECK: while.body.i34.i: // CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48 -// CHECK-NEXT: [[OR_COND_I_I:%.*]] = icmp ult i8 [[TMP3]], 10 -// CHECK-NEXT: br i1 [[OR_COND_I_I]], label [[IF_END31_I_I:%.*]], label [[IF_ELSE_I_I:%.*]] +// CHECK-NEXT: [[OR_COND_I35_I:%.*]] = icmp ult i8 [[TMP3]], 10 +// CHECK-NEXT: br i1 [[OR_COND_I35_I]], label [[IF_END31_I_I:%.*]], label [[IF_ELSE_I_I:%.*]] // CHECK: if.else.i.i: // CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97 // CHECK-NEXT: [[OR_COND33_I_I:%.*]] = icmp ult i8 [[TMP4]], 6 @@ -165,33 +165,54 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) { // CHECK: if.else17.i.i: // CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65 // CHECK-NEXT: [[OR_COND34_I_I:%.*]] = icmp ult i8 [[TMP5]], 6 -// CHECK-NEXT: br i1 [[OR_COND34_I_I]], label [[IF_END31_I_I]], label [[CLEANUP_I_I]] +// CHECK-NEXT: br i1 [[OR_COND34_I_I]], label [[IF_END31_I_I]], label [[CLEANUP_I36_I]] // CHECK: if.end31.i.i: -// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I]] ], [ -87, [[IF_ELSE_I_I]] ], [ -55, [[IF_ELSE17_I_I]] ] -// CHECK-NEXT: [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I_I]], 4 +// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I34_I]] ], [ -87, [[IF_ELSE_I_I]] ], [ -55, [[IF_ELSE17_I_I]] ] +// CHECK-NEXT: [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I32_I]], 4 // CHECK-NEXT: [[CONV25_I_I:%.*]] = sext i8 [[TMP2]] to i64 // CHECK-NEXT: [[ADD26_I_I:%.*]] = add i64 [[MUL24_I_I]], [[DOTSINK]] // CHECK-NEXT: [[ADD28_I_I:%.*]] = add i64 [[ADD26_I_I]], [[CONV25_I_I]] +// CHECK-NEXT: [[INCDEC_PTR_I40_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I31_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I36_I]] +// CHECK: cleanup.i36.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I37_I]] = phi ptr [ [[INCDEC_PTR_I40_I]], [[IF_END31_I_I]] ], [ [[__TAGP_ADDR_0_I31_I]], [[IF_ELSE17_I_I]] ] +// CHECK-NEXT: [[__R_2_I_I]] = phi i64 [ [[ADD28_I_I]], [[IF_END31_I_I]] ], [ [[__R_0_I32_I]], [[IF_ELSE17_I_I]] ] +// CHECK-NEXT: [[COND_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I]] ], [ false, [[IF_ELSE17_I_I]] ] +// CHECK-NEXT: br i1 [[COND_I_I]], label [[WHILE_COND_I30_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP11]] +// CHECK: while.cond.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I:%.*]], [[CLEANUP_I_I:%.*]] ], [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ] +// CHECK-NEXT: [[__R_0_I_I:%.*]] = phi i64 [ [[__R_1_I_I:%.*]], [[CLEANUP_I_I]] ], [ 0, [[IF_THEN_I]] ] +// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I]], align 1, !tbaa [[TBAA4]] +// CHECK-NEXT: [[CMP_NOT_I_I:%.*]] = icmp eq i8 [[TMP6]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I_I:%.*]] +// CHECK: while.body.i.i: +// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8 +// CHECK-NEXT: [[OR_COND_I_I:%.*]] = icmp eq i8 [[TMP7]], 48 +// CHECK-NEXT: br i1 [[OR_COND_I_I]], label [[IF_THEN_I_I:%.*]], label [[CLEANUP_I_I]] +// CHECK: if.then.i.i: +// CHECK-NEXT: [[MUL_I_I:%.*]] = shl i64 [[__R_0_I_I]], 3 +// CHECK-NEXT: [[CONV5_I_I:%.*]] = sext i8 [[TMP6]] to i64 +// CHECK-NEXT: [[ADD_I_I:%.*]] = add i64 [[MUL_I_I]], -48 +// CHECK-NEXT: [[SUB_I_I:%.*]] = add i64 [[ADD_I_I]], [[CONV5_I_I]] // CHECK-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I]], i64 1 // CHECK-NEXT: br label [[CLEANUP_I_I]] // CHECK: cleanup.i.i: -// CHECK-NEXT: [[__TAGP_ADDR_1_I_I]] = phi ptr [ [[INCDEC_PTR_I_I]], [[IF_END31_I_I]] ], [ [[__TAGP_ADDR_0_I_I]], [[IF_ELSE17_I_I]] ] -// CHECK-NEXT: [[__R_2_I_I]] = phi i64 [ [[ADD28_I_I]], [[IF_END31_I_I]] ], [ [[__R_0_I_I]], [[IF_ELSE17_I_I]] ] -// CHECK-NEXT: [[COND_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I]] ], [ false, [[IF_ELSE17_I_I]] ] -// CHECK-NEXT: br i1 [[COND_I_I]], label [[WHILE_COND_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP11]] +// CHECK-NEXT: [[__TAGP_ADDR_1_I_I]] = phi ptr [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ], [ [[__TAGP_ADDR_0_I_I]], [[WHILE_BODY_I_I]] ] +// CHECK-NEXT: [[__R_1_I_I]] = phi i64 [ [[SUB_I_I]], [[IF_THEN_I_I]] ], [ [[__R_0_I_I]], [[WHILE_BODY_I_I]] ] +// CHECK-NEXT: br i1 [[OR_COND_I_I]], label [[WHILE_COND_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP7]] // CHECK: while.cond.i14.i: -// CHECK-NEXT: [[__TAGP_ADDR_0_I15_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I:%.*]], [[CLEANUP_I20_I:%.*]] ], [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ] -// CHECK-NEXT: [[__R_0_I16_I:%.*]] = phi i64 [ [[__R_1_I22_I:%.*]], [[CLEANUP_I20_I]] ], [ 0, [[IF_THEN_I]] ] -// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I]], align 1, !tbaa [[TBAA4]] -// CHECK-NEXT: [[CMP_NOT_I17_I:%.*]] = icmp eq i8 [[TMP6]], 0 +// CHECK-NEXT: [[__TAGP_ADDR_0_I15_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I:%.*]], [[CLEANUP_I20_I:%.*]] ], [ [[P]], [[ENTRY:%.*]] ] +// CHECK-NEXT: [[__R_0_I16_I:%.*]] = phi i64 [ [[__R_1_I22_I:%.*]], [[CLEANUP_I20_I]] ], [ 0, [[ENTRY]] ] +// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I]], align 1, !tbaa [[TBAA4]] +// CHECK-NEXT: [[CMP_NOT_I17_I:%.*]] = icmp eq i8 [[TMP8]], 0 // CHECK-NEXT: br i1 [[CMP_NOT_I17_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I18_I:%.*]] // CHECK: while.body.i18.i: -// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8 -// CHECK-NEXT: [[OR_COND_I19_I:%.*]] = icmp eq i8 [[TMP7]], 48 +// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48 +// CHECK-NEXT: [[OR_COND_I19_I:%.*]] = icmp ult i8 [[TMP9]], 10 // CHECK-NEXT: br i1 [[OR_COND_I19_I]], label [[IF_THEN_I24_I:%.*]], label [[CLEANUP_I20_I]] // CHECK: if.then.i24.i: -// CHECK-NEXT: [[MUL_I25_I:%.*]] = shl i64 [[__R_0_I16_I]], 3 -// CHECK-NEXT: [[CONV5_I26_I:%.*]] = sext i8 [[TMP6]] to i64 +// CHECK-NEXT: [[MUL_I25_I:%.*]] = mul i64 [[__R_0_I16_I]], 10 +// CHECK-NEXT: [[CONV5_I26_I:%.*]] = sext i8 [[TMP8]] to i64 // CHECK-NEXT: [[ADD_I27_I:%.*]] = add i64 [[MUL_I25_I]], -48 // CHECK-NEXT: [[SUB_I28_I:%.*]] = add i64 [[ADD_I27_I]], [[CONV5_I26_I]] // CHECK-NEXT: [[INCDEC_PTR_I29_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I15_I]], i64 1 @@ -199,30 +220,9 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) { // CHECK: cleanup.i20.i: // CHECK-NEXT: [[__TAGP_ADDR_1_I21_I]] = phi ptr [ [[INCDEC_PTR_I29_I]], [[IF_THEN_I24_I]] ], [ [[__TAGP_ADDR_0_I15_I]], [[WHILE_BODY_I18_I]] ] // CHECK-NEXT: [[__R_1_I22_I]] = phi i64 [ [[SUB_I28_I]], [[IF_THEN_I24_I]] ], [ [[__R_0_I16_I]], [[WHILE_BODY_I18_I]] ] -// CHECK-NEXT: br i1 [[OR_COND_I19_I]], label [[WHILE_COND_I14_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP7]] -// CHECK: while.cond.i30.i: -// CHECK-NEXT: [[__TAGP_ADDR_0_I31_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I:%.*]], [[CLEANUP_I36_I:%.*]] ], [ [[P]], [[ENTRY:%.*]] ] -// CHECK-NEXT: [[__R_0_I32_I:%.*]] = phi i64 [ [[__R_1_I38_I:%.*]], [[CLEANUP_I36_I]] ], [ 0, [[ENTRY]] ] -// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I]], align 1, !tbaa [[TBAA4]] -// CHECK-NEXT: [[CMP_NOT_I33_I:%.*]] = icmp eq i8 [[TMP8]], 0 -// CHECK-NEXT: br i1 [[CMP_NOT_I33_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I34_I:%.*]] -// CHECK: while.body.i34.i: -// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48 -// CHECK-NEXT: [[OR_COND_I35_I:%.*]] = icmp ult i8 [[TMP9]], 10 -// CHECK-NEXT: br i1 [[OR_COND_I35_I]], label [[IF_THEN_I40_I:%.*]], label [[CLEANUP_I36_I]] -// CHECK: if.then.i40.i: -// CHECK-NEXT: [[MUL_I41_I:%.*]] = mul i64 [[__R_0_I32_I]], 10 -// CHECK-NEXT: [[CONV5_I42_I:%.*]] = sext i8 [[TMP8]] to i64 -// CHECK-NEXT: [[ADD_I43_I:%.*]] = add i64 [[MUL_I41_I]], -48 -// CHECK-NEXT: [[SUB_I44_I:%.*]] = add i64 [[ADD_I43_I]], [[CONV5_I42_I]] -// CHECK-NEXT: [[INCDEC_PTR_I45_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I31_I]], i64 1 -// CHECK-NEXT: br label [[CLEANUP_I36_I]] -// CHECK: cleanup.i36.i: -// CHECK-NEXT: [[__TAGP_ADDR_1_I37_I]] = phi ptr [ [[INCDEC_PTR_I45_I]], [[IF_THEN_I40_I]] ], [ [[__TAGP_ADDR_0_I31_I]], [[WHILE_BODY_I34_I]] ] -// CHECK-NEXT: [[__R_1_I38_I]] = phi i64 [ [[SUB_I44_I]], [[IF_THEN_I40_I]] ], [ [[__R_0_I32_I]], [[WHILE_BODY_I34_I]] ] -// CHECK-NEXT: br i1 [[OR_COND_I35_I]], label [[WHILE_COND_I30_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP10]] +// CHECK-NEXT: br i1 [[OR_COND_I19_I]], label [[WHILE_COND_I14_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP10]] // CHECK: _ZL15__make_mantissaPKc.exit: -// CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, [[CLEANUP_I20_I]] ], [ [[__R_0_I16_I]], [[WHILE_COND_I14_I]] ], [ 0, [[CLEANUP_I_I]] ], [ [[__R_0_I_I]], [[WHILE_COND_I_I]] ], [ 0, [[CLEANUP_I36_I]] ], [ [[__R_0_I32_I]], [[WHILE_COND_I30_I]] ] +// CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I]] ], [ [[__R_0_I_I]], [[WHILE_COND_I_I]] ], [ 0, [[CLEANUP_I36_I]] ], [ [[__R_0_I32_I]], [[WHILE_COND_I30_I]] ], [ 0, [[CLEANUP_I20_I]] ], [ [[__R_0_I16_I]], [[WHILE_COND_I14_I]] ] // CHECK-NEXT: ret i64 [[RETVAL_0_I]] // extern "C" __device__ uint64_t test___make_mantissa(const char *p) { @@ -231,8 +231,8 @@ extern "C" __device__ uint64_t test___make_mantissa(const char *p) { // CHECK-LABEL: @test_abs( // CHECK-NEXT: entry: -// CHECK-NEXT: [[ABS_I:%.*]] = tail call noundef i32 @llvm.abs.i32(i32 [[X:%.*]], i1 true) -// CHECK-NEXT: ret i32 [[ABS_I]] +// CHECK-NEXT: [[TMP0:%.*]] = tail call noundef i32 @llvm.abs.i32(i32 [[X:%.*]], i1 true) +// CHECK-NEXT: ret i32 [[TMP0]] // extern "C" __device__ int test_abs(int x) { return abs(x); @@ -240,8 +240,8 @@ extern "C" __device__ int test_abs(int x) { // CHECK-LABEL: @test_labs( // CHECK-NEXT: entry: -// CHECK-NEXT: [[ABS_I:%.*]] = tail call noundef i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true) -// CHECK-NEXT: ret i64 [[ABS_I]] +// CHECK-NEXT: [[TMP0:%.*]] = tail call noundef i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true) +// CHECK-NEXT: ret i64 [[TMP0]] // extern "C" __device__ long test_labs(long x) { return labs(x); @@ -249,8 +249,8 @@ extern "C" __device__ long test_labs(long x) { // CHECK-LABEL: @test_llabs( // CHECK-NEXT: entry: -// CHECK-NEXT: [[ABS_I:%.*]] = tail call noundef i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true) -// CHECK-NEXT: ret i64 [[ABS_I]] +// CHECK-NEXT: [[TMP0:%.*]] = tail call noundef i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true) +// CHECK-NEXT: ret i64 [[TMP0]] // extern "C" __device__ long long test_llabs(long x) { return llabs(x); @@ -649,8 +649,8 @@ extern "C" __device__ double test_copysign(double x, double y) { // // APPROX-LABEL: @test_cosf( // APPROX-NEXT: entry: -// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR16:[0-9]+]] -// APPROX-NEXT: ret float [[CALL_I_I]] +// APPROX-NEXT: [[CALL_I1:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR16:[0-9]+]] +// APPROX-NEXT: ret float [[CALL_I1]] // extern "C" __device__ float test_cosf(float x) { return cosf(x); @@ -1670,30 +1670,30 @@ extern "C" __device__ double test_j1(double x) { // DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // DEFAULT-NEXT: ] // DEFAULT: if.then.i: -// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]] // DEFAULT: if.then2.i: -// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]] // DEFAULT-NEXT: br label [[_ZL3JNFIF_EXIT]] // DEFAULT: if.end4.i: -// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 -// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]] +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1 +// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]] // DEFAULT: for.body.i: -// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] -// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] -// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] -// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// DEFAULT-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1 // DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float // DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]] -// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]] -// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]] -// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 -// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] -// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]] +// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]] +// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]] +// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1 +// DEFAULT-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// DEFAULT-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]] // DEFAULT: _ZL3jnfif.exit: -// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] // DEFAULT-NEXT: ret float [[RETVAL_0_I]] // // FINITEONLY-LABEL: @test_jnf( @@ -1703,30 +1703,30 @@ extern "C" __device__ double test_j1(double x) { // FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // FINITEONLY-NEXT: ] // FINITEONLY: if.then.i: -// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j0_f32(float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j0_f32(float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]] // FINITEONLY: if.then2.i: -// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT]] // FINITEONLY: if.end4.i: -// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j0_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 -// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]] +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j0_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_j1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1 +// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]] // FINITEONLY: for.body.i: -// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] -// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] -// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] -// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// FINITEONLY-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1 // FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float // FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]] -// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_025_I]], [[DIV_I]] -// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_024_I]] -// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 -// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] -// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]] +// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_0_I3]], [[DIV_I]] +// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_0_I2]] +// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1 +// FINITEONLY-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]] // FINITEONLY: _ZL3jnfif.exit: -// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] // FINITEONLY-NEXT: ret float [[RETVAL_0_I]] // // APPROX-LABEL: @test_jnf( @@ -1736,30 +1736,30 @@ extern "C" __device__ double test_j1(double x) { // APPROX-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // APPROX-NEXT: ] // APPROX: if.then.i: -// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]] +// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]] // APPROX-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]] // APPROX: if.then2.i: -// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]] +// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]] // APPROX-NEXT: br label [[_ZL3JNFIF_EXIT]] // APPROX: if.end4.i: -// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR16]] -// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]] -// APPROX-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 -// APPROX-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]] +// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR16]] +// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]] +// APPROX-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1 +// APPROX-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]] // APPROX: for.body.i: -// APPROX-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] -// APPROX-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] -// APPROX-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] -// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// APPROX-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// APPROX-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// APPROX-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ] +// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1 // APPROX-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float // APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]] -// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]] -// APPROX-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]] -// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 -// APPROX-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] -// APPROX-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]] +// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]] +// APPROX-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]] +// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1 +// APPROX-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// APPROX-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]] // APPROX: _ZL3jnfif.exit: -// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] // APPROX-NEXT: ret float [[RETVAL_0_I]] // extern "C" __device__ float test_jnf(int x, float y) { @@ -1773,30 +1773,30 @@ extern "C" __device__ float test_jnf(int x, float y) { // DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // DEFAULT-NEXT: ] // DEFAULT: if.then.i: -// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: br label [[_ZL2JNID_EXIT:%.*]] // DEFAULT: if.then2.i: -// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]] // DEFAULT-NEXT: br label [[_ZL2JNID_EXIT]] // DEFAULT: if.end4.i: -// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 -// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]] +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1 +// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]] // DEFAULT: for.body.i: -// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] -// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] -// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] -// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// DEFAULT-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1 // DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double // DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]] -// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]] -// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]] -// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 -// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] -// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP15:![0-9]+]] +// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]] +// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]] +// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1 +// DEFAULT-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// DEFAULT-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP15:![0-9]+]] // DEFAULT: _ZL2jnid.exit: -// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] // DEFAULT-NEXT: ret double [[RETVAL_0_I]] // // FINITEONLY-LABEL: @test_jn( @@ -1806,30 +1806,30 @@ extern "C" __device__ float test_jnf(int x, float y) { // FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // FINITEONLY-NEXT: ] // FINITEONLY: if.then.i: -// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j0_f64(double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j0_f64(double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT:%.*]] // FINITEONLY: if.then2.i: -// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT]] // FINITEONLY: if.end4.i: -// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j0_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 -// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]] +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j0_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_j1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1 +// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]] // FINITEONLY: for.body.i: -// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] -// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] -// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] -// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// FINITEONLY-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1 // FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double // FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]] -// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_025_I]], [[DIV_I]] -// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_024_I]] -// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 -// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] -// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP15:![0-9]+]] +// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_0_I3]], [[DIV_I]] +// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_0_I2]] +// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1 +// FINITEONLY-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP15:![0-9]+]] // FINITEONLY: _ZL2jnid.exit: -// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] // FINITEONLY-NEXT: ret double [[RETVAL_0_I]] // // APPROX-LABEL: @test_jn( @@ -1839,30 +1839,30 @@ extern "C" __device__ float test_jnf(int x, float y) { // APPROX-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // APPROX-NEXT: ] // APPROX: if.then.i: -// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]] +// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]] // APPROX-NEXT: br label [[_ZL2JNID_EXIT:%.*]] // APPROX: if.then2.i: -// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]] +// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]] // APPROX-NEXT: br label [[_ZL2JNID_EXIT]] // APPROX: if.end4.i: -// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR16]] -// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]] -// APPROX-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 -// APPROX-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]] +// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR16]] +// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]] +// APPROX-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1 +// APPROX-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]] // APPROX: for.body.i: -// APPROX-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] -// APPROX-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] -// APPROX-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] -// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// APPROX-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// APPROX-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// APPROX-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ] +// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1 // APPROX-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double // APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]] -// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]] -// APPROX-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]] -// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 -// APPROX-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] -// APPROX-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP15:![0-9]+]] +// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]] +// APPROX-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]] +// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1 +// APPROX-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// APPROX-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP15:![0-9]+]] // APPROX: _ZL2jnid.exit: -// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] // APPROX-NEXT: ret double [[RETVAL_0_I]] // extern "C" __device__ double test_jn(int x, double y) { @@ -2364,26 +2364,26 @@ extern "C" __device__ double test_modf(double x, double* y) { // CHECK-NEXT: entry: // CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa [[TBAA4]] // CHECK-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48 -// CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I30_I_I:%.*]] +// CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I14_I_I:%.*]] // CHECK: if.then.i.i: // CHECK-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[TAG]], i64 1 // CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[TBAA4]] -// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I14_I_I:%.*]] [ -// CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_I_PREHEADER:%.*]] -// CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_I_PREHEADER]] +// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [ +// CHECK-NEXT: i8 120, label [[WHILE_COND_I30_I_I_PREHEADER:%.*]] +// CHECK-NEXT: i8 88, label [[WHILE_COND_I30_I_I_PREHEADER]] // CHECK-NEXT: ] -// CHECK: while.cond.i.i.i.preheader: -// CHECK-NEXT: br label [[WHILE_COND_I_I_I:%.*]] -// CHECK: while.cond.i.i.i: -// CHECK-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I_I_I_PREHEADER]] ] -// CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[WHILE_COND_I_I_I_PREHEADER]] ] -// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA4]] -// CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP2]], 0 -// CHECK-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL4NANFPKC_EXIT:%.*]], label [[WHILE_BODY_I_I_I:%.*]] -// CHECK: while.body.i.i.i: +// CHECK: while.cond.i30.i.i.preheader: +// CHECK-NEXT: br label [[WHILE_COND_I30_I_I:%.*]] +// CHECK: while.cond.i30.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I_I:%.*]], [[CLEANUP_I36_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I30_I_I_PREHEADER]] ] +// CHECK-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I36_I_I]] ], [ 0, [[WHILE_COND_I30_I_I_PREHEADER]] ] +// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align 1, !tbaa [[TBAA4]] +// CHECK-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I33_I_I]], label [[_ZL4NANFPKC_EXIT:%.*]], label [[WHILE_BODY_I34_I_I:%.*]] +// CHECK: while.body.i34.i.i: // CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48 -// CHECK-NEXT: [[OR_COND_I_I_I:%.*]] = icmp ult i8 [[TMP3]], 10 -// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]] +// CHECK-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10 +// CHECK-NEXT: br i1 [[OR_COND_I35_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]] // CHECK: if.else.i.i.i: // CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97 // CHECK-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6 @@ -2391,33 +2391,54 @@ extern "C" __device__ double test_modf(double x, double* y) { // CHECK: if.else17.i.i.i: // CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65 // CHECK-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6 -// CHECK-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I_I_I]] +// CHECK-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I36_I_I]] // CHECK: if.end31.i.i.i: -// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ] -// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 4 +// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I34_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4 // CHECK-NEXT: [[CONV25_I_I_I:%.*]] = sext i8 [[TMP2]] to i64 // CHECK-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]] // CHECK-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]] +// CHECK-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I31_I_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I36_I_I]] +// CHECK: cleanup.i36.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]], [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I30_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP11]] +// CHECK: while.cond.i.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ] +// CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[IF_THEN_I_I]] ] +// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA4]] +// CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I_I_I:%.*]] +// CHECK: while.body.i.i.i: +// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8 +// CHECK-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48 +// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_THEN_I_I_I:%.*]], label [[CLEANUP_I_I_I]] +// CHECK: if.then.i.i.i: +// CHECK-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3 +// CHECK-NEXT: [[CONV5_I_I_I:%.*]] = sext i8 [[TMP6]] to i64 +// CHECK-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48 +// CHECK-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]] // CHECK-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1 // CHECK-NEXT: br label [[CLEANUP_I_I_I]] // CHECK: cleanup.i.i.i: -// CHECK-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[IF_ELSE17_I_I_I]] ] -// CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I_I_I]], [[IF_ELSE17_I_I_I]] ] -// CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ] -// CHECK-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP11]] +// CHECK-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[WHILE_BODY_I_I_I]] ] +// CHECK-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_BODY_I_I_I]] ] +// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP7]] // CHECK: while.cond.i14.i.i: -// CHECK-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I_I:%.*]], [[CLEANUP_I20_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ] -// CHECK-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]], [[CLEANUP_I20_I_I]] ], [ 0, [[IF_THEN_I_I]] ] -// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[TBAA4]] -// CHECK-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP6]], 0 +// CHECK-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I_I:%.*]], [[CLEANUP_I20_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ] +// CHECK-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]], [[CLEANUP_I20_I_I]] ], [ 0, [[ENTRY]] ] +// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[TBAA4]] +// CHECK-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0 // CHECK-NEXT: br i1 [[CMP_NOT_I17_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I18_I_I:%.*]] // CHECK: while.body.i18.i.i: -// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8 -// CHECK-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp eq i8 [[TMP7]], 48 +// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48 +// CHECK-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10 // CHECK-NEXT: br i1 [[OR_COND_I19_I_I]], label [[IF_THEN_I24_I_I:%.*]], label [[CLEANUP_I20_I_I]] // CHECK: if.then.i24.i.i: -// CHECK-NEXT: [[MUL_I25_I_I:%.*]] = shl i64 [[__R_0_I16_I_I]], 3 -// CHECK-NEXT: [[CONV5_I26_I_I:%.*]] = sext i8 [[TMP6]] to i64 +// CHECK-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10 +// CHECK-NEXT: [[CONV5_I26_I_I:%.*]] = sext i8 [[TMP8]] to i64 // CHECK-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48 // CHECK-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]], [[CONV5_I26_I_I]] // CHECK-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I15_I_I]], i64 1 @@ -2425,30 +2446,9 @@ extern "C" __device__ double test_modf(double x, double* y) { // CHECK: cleanup.i20.i.i: // CHECK-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [ [[INCDEC_PTR_I29_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]], [[WHILE_BODY_I18_I_I]] ] // CHECK-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_BODY_I18_I_I]] ] -// CHECK-NEXT: br i1 [[OR_COND_I19_I_I]], label [[WHILE_COND_I14_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP7]] -// CHECK: while.cond.i30.i.i: -// CHECK-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I_I:%.*]], [[CLEANUP_I36_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ] -// CHECK-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_1_I38_I_I:%.*]], [[CLEANUP_I36_I_I]] ], [ 0, [[ENTRY]] ] -// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align 1, !tbaa [[TBAA4]] -// CHECK-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP8]], 0 -// CHECK-NEXT: br i1 [[CMP_NOT_I33_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I34_I_I:%.*]] -// CHECK: while.body.i34.i.i: -// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48 -// CHECK-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP9]], 10 -// CHECK-NEXT: br i1 [[OR_COND_I35_I_I]], label [[IF_THEN_I40_I_I:%.*]], label [[CLEANUP_I36_I_I]] -// CHECK: if.then.i40.i.i: -// CHECK-NEXT: [[MUL_I41_I_I:%.*]] = mul i64 [[__R_0_I32_I_I]], 10 -// CHECK-NEXT: [[CONV5_I42_I_I:%.*]] = sext i8 [[TMP8]] to i64 -// CHECK-NEXT: [[ADD_I43_I_I:%.*]] = add i64 [[MUL_I41_I_I]], -48 -// CHECK-NEXT: [[SUB_I44_I_I:%.*]] = add i64 [[ADD_I43_I_I]], [[CONV5_I42_I_I]] -// CHECK-NEXT: [[INCDEC_PTR_I45_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I31_I_I]], i64 1 -// CHECK-NEXT: br label [[CLEANUP_I36_I_I]] -// CHECK: cleanup.i36.i.i: -// CHECK-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [ [[INCDEC_PTR_I45_I_I]], [[IF_THEN_I40_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]], [[WHILE_BODY_I34_I_I]] ] -// CHECK-NEXT: [[__R_1_I38_I_I]] = phi i64 [ [[SUB_I44_I_I]], [[IF_THEN_I40_I_I]] ], [ [[__R_0_I32_I_I]], [[WHILE_BODY_I34_I_I]] ] -// CHECK-NEXT: br i1 [[OR_COND_I35_I_I]], label [[WHILE_COND_I30_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP10]] +// CHECK-NEXT: br i1 [[OR_COND_I19_I_I]], label [[WHILE_COND_I14_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP10]] // CHECK: _ZL4nanfPKc.exit: -// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I20_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_COND_I14_I_I]] ], [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I36_I_I]] ], [ [[__R_0_I32_I_I]], [[WHILE_COND_I30_I_I]] ] +// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I36_I_I]] ], [ [[__R_0_I32_I_I]], [[WHILE_COND_I30_I_I]] ], [ 0, [[CLEANUP_I20_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_COND_I14_I_I]] ] // CHECK-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32 // CHECK-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303 // CHECK-NEXT: [[BF_SET9_I:%.*]] = or i32 [[BF_VALUE_I]], 2143289344 @@ -2463,26 +2463,26 @@ extern "C" __device__ float test_nanf(const char *tag) { // CHECK-NEXT: entry: // CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa [[TBAA4]] // CHECK-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48 -// CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I30_I_I:%.*]] +// CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I14_I_I:%.*]] // CHECK: if.then.i.i: // CHECK-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[TAG]], i64 1 // CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[TBAA4]] -// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I14_I_I:%.*]] [ -// CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_I_PREHEADER:%.*]] -// CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_I_PREHEADER]] +// CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I_I_I:%.*]] [ +// CHECK-NEXT: i8 120, label [[WHILE_COND_I30_I_I_PREHEADER:%.*]] +// CHECK-NEXT: i8 88, label [[WHILE_COND_I30_I_I_PREHEADER]] // CHECK-NEXT: ] -// CHECK: while.cond.i.i.i.preheader: -// CHECK-NEXT: br label [[WHILE_COND_I_I_I:%.*]] -// CHECK: while.cond.i.i.i: -// CHECK-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I_I_I_PREHEADER]] ] -// CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[WHILE_COND_I_I_I_PREHEADER]] ] -// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA4]] -// CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP2]], 0 -// CHECK-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL3NANPKC_EXIT:%.*]], label [[WHILE_BODY_I_I_I:%.*]] -// CHECK: while.body.i.i.i: +// CHECK: while.cond.i30.i.i.preheader: +// CHECK-NEXT: br label [[WHILE_COND_I30_I_I:%.*]] +// CHECK: while.cond.i30.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I_I:%.*]], [[CLEANUP_I36_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I30_I_I_PREHEADER]] ] +// CHECK-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I36_I_I]] ], [ 0, [[WHILE_COND_I30_I_I_PREHEADER]] ] +// CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align 1, !tbaa [[TBAA4]] +// CHECK-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP2]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I33_I_I]], label [[_ZL3NANPKC_EXIT:%.*]], label [[WHILE_BODY_I34_I_I:%.*]] +// CHECK: while.body.i34.i.i: // CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48 -// CHECK-NEXT: [[OR_COND_I_I_I:%.*]] = icmp ult i8 [[TMP3]], 10 -// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]] +// CHECK-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP3]], 10 +// CHECK-NEXT: br i1 [[OR_COND_I35_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]] // CHECK: if.else.i.i.i: // CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97 // CHECK-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6 @@ -2490,33 +2490,54 @@ extern "C" __device__ float test_nanf(const char *tag) { // CHECK: if.else17.i.i.i: // CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65 // CHECK-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6 -// CHECK-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I_I_I]] +// CHECK-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I36_I_I]] // CHECK: if.end31.i.i.i: -// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ] -// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 4 +// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I34_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I32_I_I]], 4 // CHECK-NEXT: [[CONV25_I_I_I:%.*]] = sext i8 [[TMP2]] to i64 // CHECK-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]] // CHECK-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]] +// CHECK-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I31_I_I]], i64 1 +// CHECK-NEXT: br label [[CLEANUP_I36_I_I]] +// CHECK: cleanup.i36.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]], [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I32_I_I]], [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ] +// CHECK-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I30_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP11]] +// CHECK: while.cond.i.i.i: +// CHECK-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ] +// CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_1_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[IF_THEN_I_I]] ] +// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA4]] +// CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP6]], 0 +// CHECK-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I_I_I:%.*]] +// CHECK: while.body.i.i.i: +// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8 +// CHECK-NEXT: [[OR_COND_I_I_I:%.*]] = icmp eq i8 [[TMP7]], 48 +// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_THEN_I_I_I:%.*]], label [[CLEANUP_I_I_I]] +// CHECK: if.then.i.i.i: +// CHECK-NEXT: [[MUL_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 3 +// CHECK-NEXT: [[CONV5_I_I_I:%.*]] = sext i8 [[TMP6]] to i64 +// CHECK-NEXT: [[ADD_I_I_I:%.*]] = add i64 [[MUL_I_I_I]], -48 +// CHECK-NEXT: [[SUB_I_I_I:%.*]] = add i64 [[ADD_I_I_I]], [[CONV5_I_I_I]] // CHECK-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1 // CHECK-NEXT: br label [[CLEANUP_I_I_I]] // CHECK: cleanup.i.i.i: -// CHECK-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[IF_ELSE17_I_I_I]] ] -// CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I_I_I]], [[IF_ELSE17_I_I_I]] ] -// CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ] -// CHECK-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP11]] +// CHECK-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[WHILE_BODY_I_I_I]] ] +// CHECK-NEXT: [[__R_1_I_I_I]] = phi i64 [ [[SUB_I_I_I]], [[IF_THEN_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_BODY_I_I_I]] ] +// CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP7]] // CHECK: while.cond.i14.i.i: -// CHECK-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I_I:%.*]], [[CLEANUP_I20_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ] -// CHECK-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]], [[CLEANUP_I20_I_I]] ], [ 0, [[IF_THEN_I_I]] ] -// CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[TBAA4]] -// CHECK-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP6]], 0 +// CHECK-NEXT: [[__TAGP_ADDR_0_I15_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I21_I_I:%.*]], [[CLEANUP_I20_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ] +// CHECK-NEXT: [[__R_0_I16_I_I:%.*]] = phi i64 [ [[__R_1_I22_I_I:%.*]], [[CLEANUP_I20_I_I]] ], [ 0, [[ENTRY]] ] +// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I15_I_I]], align 1, !tbaa [[TBAA4]] +// CHECK-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP8]], 0 // CHECK-NEXT: br i1 [[CMP_NOT_I17_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I18_I_I:%.*]] // CHECK: while.body.i18.i.i: -// CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8 -// CHECK-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp eq i8 [[TMP7]], 48 +// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48 +// CHECK-NEXT: [[OR_COND_I19_I_I:%.*]] = icmp ult i8 [[TMP9]], 10 // CHECK-NEXT: br i1 [[OR_COND_I19_I_I]], label [[IF_THEN_I24_I_I:%.*]], label [[CLEANUP_I20_I_I]] // CHECK: if.then.i24.i.i: -// CHECK-NEXT: [[MUL_I25_I_I:%.*]] = shl i64 [[__R_0_I16_I_I]], 3 -// CHECK-NEXT: [[CONV5_I26_I_I:%.*]] = sext i8 [[TMP6]] to i64 +// CHECK-NEXT: [[MUL_I25_I_I:%.*]] = mul i64 [[__R_0_I16_I_I]], 10 +// CHECK-NEXT: [[CONV5_I26_I_I:%.*]] = sext i8 [[TMP8]] to i64 // CHECK-NEXT: [[ADD_I27_I_I:%.*]] = add i64 [[MUL_I25_I_I]], -48 // CHECK-NEXT: [[SUB_I28_I_I:%.*]] = add i64 [[ADD_I27_I_I]], [[CONV5_I26_I_I]] // CHECK-NEXT: [[INCDEC_PTR_I29_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I15_I_I]], i64 1 @@ -2524,30 +2545,9 @@ extern "C" __device__ float test_nanf(const char *tag) { // CHECK: cleanup.i20.i.i: // CHECK-NEXT: [[__TAGP_ADDR_1_I21_I_I]] = phi ptr [ [[INCDEC_PTR_I29_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__TAGP_ADDR_0_I15_I_I]], [[WHILE_BODY_I18_I_I]] ] // CHECK-NEXT: [[__R_1_I22_I_I]] = phi i64 [ [[SUB_I28_I_I]], [[IF_THEN_I24_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_BODY_I18_I_I]] ] -// CHECK-NEXT: br i1 [[OR_COND_I19_I_I]], label [[WHILE_COND_I14_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP7]] -// CHECK: while.cond.i30.i.i: -// CHECK-NEXT: [[__TAGP_ADDR_0_I31_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I37_I_I:%.*]], [[CLEANUP_I36_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ] -// CHECK-NEXT: [[__R_0_I32_I_I:%.*]] = phi i64 [ [[__R_1_I38_I_I:%.*]], [[CLEANUP_I36_I_I]] ], [ 0, [[ENTRY]] ] -// CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I31_I_I]], align 1, !tbaa [[TBAA4]] -// CHECK-NEXT: [[CMP_NOT_I33_I_I:%.*]] = icmp eq i8 [[TMP8]], 0 -// CHECK-NEXT: br i1 [[CMP_NOT_I33_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I34_I_I:%.*]] -// CHECK: while.body.i34.i.i: -// CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48 -// CHECK-NEXT: [[OR_COND_I35_I_I:%.*]] = icmp ult i8 [[TMP9]], 10 -// CHECK-NEXT: br i1 [[OR_COND_I35_I_I]], label [[IF_THEN_I40_I_I:%.*]], label [[CLEANUP_I36_I_I]] -// CHECK: if.then.i40.i.i: -// CHECK-NEXT: [[MUL_I41_I_I:%.*]] = mul i64 [[__R_0_I32_I_I]], 10 -// CHECK-NEXT: [[CONV5_I42_I_I:%.*]] = sext i8 [[TMP8]] to i64 -// CHECK-NEXT: [[ADD_I43_I_I:%.*]] = add i64 [[MUL_I41_I_I]], -48 -// CHECK-NEXT: [[SUB_I44_I_I:%.*]] = add i64 [[ADD_I43_I_I]], [[CONV5_I42_I_I]] -// CHECK-NEXT: [[INCDEC_PTR_I45_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I31_I_I]], i64 1 -// CHECK-NEXT: br label [[CLEANUP_I36_I_I]] -// CHECK: cleanup.i36.i.i: -// CHECK-NEXT: [[__TAGP_ADDR_1_I37_I_I]] = phi ptr [ [[INCDEC_PTR_I45_I_I]], [[IF_THEN_I40_I_I]] ], [ [[__TAGP_ADDR_0_I31_I_I]], [[WHILE_BODY_I34_I_I]] ] -// CHECK-NEXT: [[__R_1_I38_I_I]] = phi i64 [ [[SUB_I44_I_I]], [[IF_THEN_I40_I_I]] ], [ [[__R_0_I32_I_I]], [[WHILE_BODY_I34_I_I]] ] -// CHECK-NEXT: br i1 [[OR_COND_I35_I_I]], label [[WHILE_COND_I30_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP10]] +// CHECK-NEXT: br i1 [[OR_COND_I19_I_I]], label [[WHILE_COND_I14_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP10]] // CHECK: _ZL3nanPKc.exit: -// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I20_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_COND_I14_I_I]] ], [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I36_I_I]] ], [ [[__R_0_I32_I_I]], [[WHILE_COND_I30_I_I]] ] +// CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I36_I_I]] ], [ [[__R_0_I32_I_I]], [[WHILE_COND_I30_I_I]] ], [ 0, [[CLEANUP_I20_I_I]] ], [ [[__R_0_I16_I_I]], [[WHILE_COND_I14_I_I]] ] // CHECK-NEXT: [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]], 2251799813685247 // CHECK-NEXT: [[BF_SET9_I:%.*]] = or i64 [[BF_VALUE_I]], 9221120237041090560 // CHECK-NEXT: [[TMP10:%.*]] = bitcast i64 [[BF_SET9_I]] to double @@ -2819,62 +2819,62 @@ extern "C" __device__ double test_normcdfinv(double x) { // DEFAULT-LABEL: @test_normf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 -// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// DEFAULT-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0 +// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] // DEFAULT: while.body.i: -// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] -// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] -// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] -// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 -// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA16]] +// DEFAULT-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// DEFAULT-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// DEFAULT-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1 +// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA16]] // DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]] -// DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]] -// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1 +// DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]] +// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1 // DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 // DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]] // DEFAULT: _ZL5normfiPKf.exit: -// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[__R_0_LCSSA_I]]) +// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[__R_0_I_LCSSA]]) // DEFAULT-NEXT: ret float [[TMP1]] // // FINITEONLY-LABEL: @test_normf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 -// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// FINITEONLY-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0 +// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] // FINITEONLY: while.body.i: -// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] -// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] -// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] -// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 -// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA16]] +// FINITEONLY-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// FINITEONLY-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1 +// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA16]] // FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]] -// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_08_I]], [[MUL_I]] -// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1 +// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_0_I4]], [[MUL_I]] +// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1 // FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]] // FINITEONLY: _ZL5normfiPKf.exit: -// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract noundef float @llvm.sqrt.f32(float [[__R_0_LCSSA_I]]) +// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract noundef float @llvm.sqrt.f32(float [[__R_0_I_LCSSA]]) // FINITEONLY-NEXT: ret float [[TMP1]] // // APPROX-LABEL: @test_normf( // APPROX-NEXT: entry: -// APPROX-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 -// APPROX-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// APPROX-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0 +// APPROX-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] // APPROX: while.body.i: -// APPROX-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] -// APPROX-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] -// APPROX-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] -// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 -// APPROX-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA16]] +// APPROX-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// APPROX-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// APPROX-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1 +// APPROX-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA16]] // APPROX-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]] -// APPROX-NEXT: [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]] -// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1 +// APPROX-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]] +// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1 // APPROX-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 // APPROX-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]] // APPROX: _ZL5normfiPKf.exit: -// APPROX-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[__R_0_LCSSA_I]]) +// APPROX-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[__R_0_I_LCSSA]]) // APPROX-NEXT: ret float [[TMP1]] // extern "C" __device__ float test_normf(int x, const float *y) { @@ -2883,62 +2883,62 @@ extern "C" __device__ float test_normf(int x, const float *y) { // DEFAULT-LABEL: @test_norm( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 -// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// DEFAULT-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0 +// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] // DEFAULT: while.body.i: -// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] -// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] -// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] -// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 -// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA18]] +// DEFAULT-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// DEFAULT-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// DEFAULT-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1 +// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA18]] // DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]] -// DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]] -// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1 +// DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]] +// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1 // DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 // DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]] // DEFAULT: _ZL4normiPKd.exit: -// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[__R_0_LCSSA_I]]) +// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[__R_0_I_LCSSA]]) // DEFAULT-NEXT: ret double [[TMP1]] // // FINITEONLY-LABEL: @test_norm( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 -// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// FINITEONLY-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0 +// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] // FINITEONLY: while.body.i: -// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] -// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] -// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] -// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 -// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA18]] +// FINITEONLY-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// FINITEONLY-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1 +// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA18]] // FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]] -// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_08_I]], [[MUL_I]] -// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1 +// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_0_I4]], [[MUL_I]] +// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1 // FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]] // FINITEONLY: _ZL4normiPKd.exit: -// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract noundef double @llvm.sqrt.f64(double [[__R_0_LCSSA_I]]) +// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract noundef double @llvm.sqrt.f64(double [[__R_0_I_LCSSA]]) // FINITEONLY-NEXT: ret double [[TMP1]] // // APPROX-LABEL: @test_norm( // APPROX-NEXT: entry: -// APPROX-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 -// APPROX-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// APPROX-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0 +// APPROX-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] // APPROX: while.body.i: -// APPROX-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] -// APPROX-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] -// APPROX-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] -// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 -// APPROX-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA18]] +// APPROX-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// APPROX-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// APPROX-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1 +// APPROX-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA18]] // APPROX-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]] -// APPROX-NEXT: [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]] -// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1 +// APPROX-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]] +// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1 // APPROX-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 // APPROX-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]] // APPROX: _ZL4normiPKd.exit: -// APPROX-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[__R_0_LCSSA_I]]) +// APPROX-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[__R_0_I_LCSSA]]) // APPROX-NEXT: ret double [[TMP1]] // extern "C" __device__ double test_norm(int x, const double *y) { @@ -3243,62 +3243,62 @@ extern "C" __device__ double test_rint(double x) { // DEFAULT-LABEL: @test_rnormf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 -// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// DEFAULT-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0 +// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] // DEFAULT: while.body.i: -// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] -// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] -// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] -// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 -// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA16]] +// DEFAULT-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// DEFAULT-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// DEFAULT-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1 +// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA16]] // DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]] -// DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]] -// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1 +// DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]] +// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1 // DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 // DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]] // DEFAULT: _ZL6rnormfiPKf.exit: -// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR15]] +// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rsqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR15]] // DEFAULT-NEXT: ret float [[CALL_I]] // // FINITEONLY-LABEL: @test_rnormf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 -// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// FINITEONLY-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0 +// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] // FINITEONLY: while.body.i: -// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] -// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] -// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] -// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 -// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA16]] +// FINITEONLY-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// FINITEONLY-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1 +// FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA16]] // FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]] -// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_08_I]], [[MUL_I]] -// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1 +// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_0_I4]], [[MUL_I]] +// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1 // FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]] // FINITEONLY: _ZL6rnormfiPKf.exit: -// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_rsqrt_f32(float noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_rsqrt_f32(float noundef nofpclass(nan inf) [[__R_0_I_LCSSA]]) #[[ATTR15]] // FINITEONLY-NEXT: ret float [[CALL_I]] // // APPROX-LABEL: @test_rnormf( // APPROX-NEXT: entry: -// APPROX-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 -// APPROX-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// APPROX-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0 +// APPROX-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] // APPROX: while.body.i: -// APPROX-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] -// APPROX-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] -// APPROX-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] -// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 -// APPROX-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA16]] +// APPROX-NEXT: [[__R_0_I4:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// APPROX-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// APPROX-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1 +// APPROX-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_0_I3]], align 4, !tbaa [[TBAA16]] // APPROX-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]] -// APPROX-NEXT: [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]] -// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1 +// APPROX-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]] +// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_0_I3]], i64 1 // APPROX-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 // APPROX-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]] // APPROX: _ZL6rnormfiPKf.exit: -// APPROX-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR15]] +// APPROX-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rsqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR15]] // APPROX-NEXT: ret float [[CALL_I]] // extern "C" __device__ float test_rnormf(int x, const float* y) { @@ -3307,62 +3307,62 @@ extern "C" __device__ float test_rnormf(int x, const float* y) { // DEFAULT-LABEL: @test_rnorm( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 -// DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// DEFAULT-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0 +// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] // DEFAULT: while.body.i: -// DEFAULT-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] -// DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] -// DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] -// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 -// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA18]] +// DEFAULT-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// DEFAULT-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// DEFAULT-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1 +// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA18]] // DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]] -// DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]] -// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1 +// DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]] +// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1 // DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 // DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]] // DEFAULT: _ZL5rnormiPKd.exit: -// DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR15]] +// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rsqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR15]] // DEFAULT-NEXT: ret double [[CALL_I]] // // FINITEONLY-LABEL: @test_rnorm( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 -// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// FINITEONLY-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0 +// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] // FINITEONLY: while.body.i: -// FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] -// FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] -// FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] -// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 -// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA18]] +// FINITEONLY-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// FINITEONLY-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1 +// FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA18]] // FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]] -// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_08_I]], [[MUL_I]] -// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1 +// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_0_I4]], [[MUL_I]] +// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1 // FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]] // FINITEONLY: _ZL5rnormiPKd.exit: -// FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_rsqrt_f64(double noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR15]] +// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_rsqrt_f64(double noundef nofpclass(nan inf) [[__R_0_I_LCSSA]]) #[[ATTR15]] // FINITEONLY-NEXT: ret double [[CALL_I]] // // APPROX-LABEL: @test_rnorm( // APPROX-NEXT: entry: -// APPROX-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0 -// APPROX-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] +// APPROX-NEXT: [[TOBOOL_NOT_I1:%.*]] = icmp eq i32 [[X:%.*]], 0 +// APPROX-NEXT: br i1 [[TOBOOL_NOT_I1]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]] // APPROX: while.body.i: -// APPROX-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] -// APPROX-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] -// APPROX-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] -// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1 -// APPROX-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA18]] +// APPROX-NEXT: [[__R_0_I4:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ] +// APPROX-NEXT: [[__A_ADDR_0_I3:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ] +// APPROX-NEXT: [[__DIM_ADDR_0_I2:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ] +// APPROX-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_0_I2]], -1 +// APPROX-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_0_I3]], align 8, !tbaa [[TBAA18]] // APPROX-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]] -// APPROX-NEXT: [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]] -// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1 +// APPROX-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]] +// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_0_I3]], i64 1 // APPROX-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0 // APPROX-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]] // APPROX: _ZL5rnormiPKd.exit: -// APPROX-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] -// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR15]] +// APPROX-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ] +// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rsqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR15]] // APPROX-NEXT: ret double [[CALL_I]] // extern "C" __device__ double test_rnorm(int x, const double* y) { @@ -3838,8 +3838,8 @@ extern "C" __device__ void test_sincospi(double x, double *y, double *z) { // // APPROX-LABEL: @test_sinf( // APPROX-NEXT: entry: -// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]] -// APPROX-NEXT: ret float [[CALL_I_I]] +// APPROX-NEXT: [[CALL_I1:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// APPROX-NEXT: ret float [[CALL_I1]] // extern "C" __device__ float test_sinf(float x) { return sinf(x); @@ -4175,30 +4175,30 @@ extern "C" __device__ double test_y1(double x) { // DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // DEFAULT-NEXT: ] // DEFAULT: if.then.i: -// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]] // DEFAULT: if.then2.i: -// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]] // DEFAULT-NEXT: br label [[_ZL3YNFIF_EXIT]] // DEFAULT: if.end4.i: -// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 -// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]] +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1 +// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]] // DEFAULT: for.body.i: -// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] -// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] -// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] -// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// DEFAULT-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1 // DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float // DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]] -// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]] -// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]] -// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 -// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] -// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]] +// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]] +// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]] +// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1 +// DEFAULT-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// DEFAULT-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]] // DEFAULT: _ZL3ynfif.exit: -// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] // DEFAULT-NEXT: ret float [[RETVAL_0_I]] // // FINITEONLY-LABEL: @test_ynf( @@ -4208,30 +4208,30 @@ extern "C" __device__ double test_y1(double x) { // FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // FINITEONLY-NEXT: ] // FINITEONLY: if.then.i: -// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y0_f32(float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y0_f32(float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]] // FINITEONLY: if.then2.i: -// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT]] // FINITEONLY: if.end4.i: -// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y0_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 -// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]] +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y0_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_y1_f32(float noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1 +// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]] // FINITEONLY: for.body.i: -// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] -// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] -// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] -// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// FINITEONLY-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1 // FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float // FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]] -// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_025_I]], [[DIV_I]] -// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_024_I]] -// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 -// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] -// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]] +// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_0_I3]], [[DIV_I]] +// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_0_I2]] +// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1 +// FINITEONLY-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]] // FINITEONLY: _ZL3ynfif.exit: -// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] // FINITEONLY-NEXT: ret float [[RETVAL_0_I]] // // APPROX-LABEL: @test_ynf( @@ -4241,30 +4241,30 @@ extern "C" __device__ double test_y1(double x) { // APPROX-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // APPROX-NEXT: ] // APPROX: if.then.i: -// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]] +// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]] // APPROX-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]] // APPROX: if.then2.i: -// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]] +// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]] // APPROX-NEXT: br label [[_ZL3YNFIF_EXIT]] // APPROX: if.end4.i: -// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR16]] -// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]] -// APPROX-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 -// APPROX-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]] +// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR16]] +// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]] +// APPROX-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1 +// APPROX-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]] // APPROX: for.body.i: -// APPROX-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] -// APPROX-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] -// APPROX-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] -// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// APPROX-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// APPROX-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// APPROX-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ] +// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1 // APPROX-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float // APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]] -// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]] -// APPROX-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]] -// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 -// APPROX-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] -// APPROX-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]] +// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]] +// APPROX-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]] +// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1 +// APPROX-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// APPROX-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]] // APPROX: _ZL3ynfif.exit: -// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] // APPROX-NEXT: ret float [[RETVAL_0_I]] // extern "C" __device__ float test_ynf(int x, float y) { @@ -4278,30 +4278,30 @@ extern "C" __device__ float test_ynf(int x, float y) { // DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // DEFAULT-NEXT: ] // DEFAULT: if.then.i: -// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]] // DEFAULT-NEXT: br label [[_ZL2YNID_EXIT:%.*]] // DEFAULT: if.then2.i: -// DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]] // DEFAULT-NEXT: br label [[_ZL2YNID_EXIT]] // DEFAULT: if.end4.i: -// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 -// DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]] +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1 +// DEFAULT-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]] // DEFAULT: for.body.i: -// DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] -// DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] -// DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] -// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// DEFAULT-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ] +// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1 // DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double // DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]] -// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]] -// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]] -// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 -// DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] -// DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]] +// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]] +// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]] +// DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1 +// DEFAULT-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// DEFAULT-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]] // DEFAULT: _ZL2ynid.exit: -// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] // DEFAULT-NEXT: ret double [[RETVAL_0_I]] // // FINITEONLY-LABEL: @test_yn( @@ -4311,30 +4311,30 @@ extern "C" __device__ float test_ynf(int x, float y) { // FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // FINITEONLY-NEXT: ] // FINITEONLY: if.then.i: -// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y0_f64(double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y0_f64(double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT:%.*]] // FINITEONLY: if.then2.i: -// FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] // FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT]] // FINITEONLY: if.end4.i: -// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y0_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 -// FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]] +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y0_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @__ocml_y1_f64(double noundef nofpclass(nan inf) [[Y]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1 +// FINITEONLY-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]] // FINITEONLY: for.body.i: -// FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] -// FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] -// FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] -// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// FINITEONLY-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ] +// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1 // FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double // FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]] -// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_025_I]], [[DIV_I]] -// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_024_I]] -// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 -// FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] -// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]] +// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_0_I3]], [[DIV_I]] +// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_0_I2]] +// FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1 +// FINITEONLY-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// FINITEONLY-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]] // FINITEONLY: _ZL2ynid.exit: -// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] // FINITEONLY-NEXT: ret double [[RETVAL_0_I]] // // APPROX-LABEL: @test_yn( @@ -4344,30 +4344,30 @@ extern "C" __device__ float test_ynf(int x, float y) { // APPROX-NEXT: i32 1, label [[IF_THEN2_I:%.*]] // APPROX-NEXT: ] // APPROX: if.then.i: -// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]] +// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]] // APPROX-NEXT: br label [[_ZL2YNID_EXIT:%.*]] // APPROX: if.then2.i: -// APPROX-NEXT: [[CALL_I20_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]] +// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]] // APPROX-NEXT: br label [[_ZL2YNID_EXIT]] // APPROX: if.end4.i: -// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR16]] -// APPROX-NEXT: [[CALL_I22_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]] -// APPROX-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1 -// APPROX-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]] +// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR16]] +// APPROX-NEXT: [[CALL_I21_I:%.*]] = tail call contract noundef double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]] +// APPROX-NEXT: [[CMP7_I1:%.*]] = icmp sgt i32 [[X]], 1 +// APPROX-NEXT: br i1 [[CMP7_I1]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]] // APPROX: for.body.i: -// APPROX-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] -// APPROX-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ] -// APPROX-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] -// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1 +// APPROX-NEXT: [[__I_0_I4:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ] +// APPROX-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ] +// APPROX-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ] +// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1 // APPROX-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double // APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]] -// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]] -// APPROX-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]] -// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1 -// APPROX-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]] -// APPROX-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]] +// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]] +// APPROX-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]] +// APPROX-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_0_I4]], 1 +// APPROX-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i32 [[INC_I]], [[X]] +// APPROX-NEXT: br i1 [[EXITCOND_NOT]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]] // APPROX: _ZL2ynid.exit: -// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] +// APPROX-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I20_I]], [[IF_THEN_I]] ], [ [[CALL_I22_I]], [[IF_THEN2_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ] // APPROX-NEXT: ret double [[RETVAL_0_I]] // extern "C" __device__ double test_yn(int x, double y) { @@ -4742,26 +4742,26 @@ extern "C" __device__ float test___sinf(float x) { // DEFAULT-LABEL: @test___tanf( // DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]] -// DEFAULT-NEXT: [[CALL_I3_I:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]] -// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.rcp.f32(float [[CALL_I3_I]]) -// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[CALL_I_I]], [[TMP0]] +// DEFAULT-NEXT: [[CALL_I3_I:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]] +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.rcp.f32(float [[CALL_I_I]]) +// DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[CALL_I3_I]], [[TMP0]] // DEFAULT-NEXT: ret float [[MUL_I]] // // FINITEONLY-LABEL: @test___tanf( // FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_native_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[CALL_I3_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_native_cos_f32(float noundef nofpclass(nan inf) [[X]]) #[[ATTR16]] -// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.amdgcn.rcp.f32(float [[CALL_I3_I]]) -// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[CALL_I_I]], [[TMP0]] +// FINITEONLY-NEXT: [[CALL_I3_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_native_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_native_cos_f32(float noundef nofpclass(nan inf) [[X]]) #[[ATTR16]] +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.amdgcn.rcp.f32(float [[CALL_I_I]]) +// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[CALL_I3_I]], [[TMP0]] // FINITEONLY-NEXT: ret float [[MUL_I]] // // APPROX-LABEL: @test___tanf( // APPROX-NEXT: entry: -// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]] -// APPROX-NEXT: [[CALL_I3_I:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]] -// APPROX-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.rcp.f32(float [[CALL_I3_I]]) -// APPROX-NEXT: [[MUL_I:%.*]] = fmul contract float [[CALL_I_I]], [[TMP0]] +// APPROX-NEXT: [[CALL_I3_I:%.*]] = tail call contract noundef float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]] +// APPROX-NEXT: [[CALL_I_I:%.*]] = tail call contract noundef float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]] +// APPROX-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.rcp.f32(float [[CALL_I_I]]) +// APPROX-NEXT: [[MUL_I:%.*]] = fmul contract float [[CALL_I3_I]], [[TMP0]] // APPROX-NEXT: ret float [[MUL_I]] // extern "C" __device__ float test___tanf(float x) { diff --git a/clang/test/OpenMP/bug57757.cpp b/clang/test/OpenMP/bug57757.cpp index 7894796ac4628..7acfe134ddd0b 100644 --- a/clang/test/OpenMP/bug57757.cpp +++ b/clang/test/OpenMP/bug57757.cpp @@ -32,24 +32,23 @@ void foo() { // CHECK-NEXT: entry: // CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [[STRUCT_KMP_TASK_T:%.*]], ptr [[TMP1]], i64 0, i32 2 // CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META13:![0-9]+]]) -// CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META16:![0-9]+]]) -// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4, !tbaa [[TBAA18:![0-9]+]], !alias.scope !13, !noalias !16 +// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4, !tbaa [[TBAA16:![0-9]+]], !alias.scope !13, !noalias !17 // CHECK-NEXT: switch i32 [[TMP3]], label [[DOTOMP_OUTLINED__EXIT:%.*]] [ // CHECK-NEXT: i32 0, label [[DOTUNTIED_JMP__I:%.*]] // CHECK-NEXT: i32 1, label [[DOTUNTIED_NEXT__I:%.*]] // CHECK-NEXT: ] // CHECK: .untied.jmp..i: -// CHECK-NEXT: store i32 1, ptr [[TMP2]], align 4, !tbaa [[TBAA18]], !alias.scope !13, !noalias !16 -// CHECK-NEXT: [[TMP4:%.*]] = tail call i32 @__kmpc_omp_task(ptr nonnull @[[GLOB1]], i32 [[TMP0]], ptr [[TMP1]]), !noalias !19 +// CHECK-NEXT: store i32 1, ptr [[TMP2]], align 4, !tbaa [[TBAA16]], !alias.scope !13, !noalias !17 +// CHECK-NEXT: [[TMP4:%.*]] = tail call i32 @__kmpc_omp_task(ptr nonnull @[[GLOB1]], i32 [[TMP0]], ptr [[TMP1]]), !noalias !13 // CHECK-NEXT: br label [[DOTOMP_OUTLINED__EXIT]] // CHECK: .untied.next..i: // CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [[STRUCT_KMP_TASK_T_WITH_PRIVATES:%.*]], ptr [[TMP1]], i64 0, i32 1 // CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [[STRUCT_KMP_TASK_T_WITH_PRIVATES]], ptr [[TMP1]], i64 0, i32 1, i32 2 // CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds [[STRUCT_KMP_TASK_T_WITH_PRIVATES]], ptr [[TMP1]], i64 0, i32 1, i32 1 -// CHECK-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP5]], align 8, !tbaa [[TBAA20:![0-9]+]], !alias.scope !16, !noalias !13 -// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP7]], align 4, !tbaa [[TBAA18]], !alias.scope !16, !noalias !13 -// CHECK-NEXT: [[TMP10:%.*]] = load float, ptr [[TMP6]], align 4, !tbaa [[TBAA21:![0-9]+]], !alias.scope !16, !noalias !13 -// CHECK-NEXT: tail call void [[TMP8]](i32 noundef [[TMP9]], float noundef [[TMP10]]) #[[ATTR2:[0-9]+]], !noalias !19 +// CHECK-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP5]], align 8, !tbaa [[TBAA19:![0-9]+]], !noalias !13 +// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP7]], align 4, !tbaa [[TBAA16]], !noalias !13 +// CHECK-NEXT: [[TMP10:%.*]] = load float, ptr [[TMP6]], align 4, !tbaa [[TBAA20:![0-9]+]], !noalias !13 +// CHECK-NEXT: tail call void [[TMP8]](i32 noundef [[TMP9]], float noundef [[TMP10]]) #[[ATTR2:[0-9]+]], !noalias !13 // CHECK-NEXT: br label [[DOTOMP_OUTLINED__EXIT]] // CHECK: .omp_outlined..exit: // CHECK-NEXT: ret i32 0 diff --git a/llvm/lib/Passes/PassBuilderPipelines.cpp b/llvm/lib/Passes/PassBuilderPipelines.cpp index 600f8d43caaf2..baea2913338cd 100644 --- a/llvm/lib/Passes/PassBuilderPipelines.cpp +++ b/llvm/lib/Passes/PassBuilderPipelines.cpp @@ -164,7 +164,7 @@ static cl::opt EnableModuleInliner("enable-module-inliner", cl::desc("Enable module inliner")); static cl::opt PerformMandatoryInliningsFirst( - "mandatory-inlining-first", cl::init(true), cl::Hidden, + "mandatory-inlining-first", cl::init(false), cl::Hidden, cl::desc("Perform mandatory inlinings module-wide, before performing " "inlining")); @@ -1127,6 +1127,8 @@ PassBuilder::buildModuleSimplificationPipeline(OptimizationLevel Level, if (EnableSyntheticCounts && !PGOOpt) MPM.addPass(SyntheticCountsPropagation()); + MPM.addPass(AlwaysInlinerPass(/*InsertLifetimeIntrinsics=*/true)); + if (EnableModuleInliner) MPM.addPass(buildModuleInlinerPipeline(Level, Phase)); else diff --git a/llvm/test/Other/new-pm-defaults.ll b/llvm/test/Other/new-pm-defaults.ll index 098db2b959a29..ecdb5a5e010d9 100644 --- a/llvm/test/Other/new-pm-defaults.ll +++ b/llvm/test/Other/new-pm-defaults.ll @@ -121,6 +121,8 @@ ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy ; CHECK-EP-PEEPHOLE-NEXT: Running pass: NoOpFunctionPass ; CHECK-O-NEXT: Running pass: SimplifyCFGPass +; CHECK-O-NEXT: Running pass: AlwaysInlinerPass +; CHECK-O-NEXT: Running analysis: ProfileSummaryAnalysis ; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass ; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA @@ -129,14 +131,12 @@ ; CHECK-O-NEXT: Running pass: InvalidateAnalysisPass<{{.*}}AAManager ; CHECK-O-NEXT: Invalidating analysis: AAManager ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}ProfileSummaryAnalysis -; CHECK-O-NEXT: Running analysis: ProfileSummaryAnalysis ; CHECK-O-NEXT: Running analysis: InnerAnalysisManagerProxy ; CHECK-O-NEXT: Running analysis: LazyCallGraphAnalysis ; CHECK-O-NEXT: Running analysis: FunctionAnalysisManagerCGSCCProxy ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy<{{.*}}LazyCallGraph::SCC{{.*}}> ; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass ; CHECK-O-NEXT: Running pass: InlinerPass -; CHECK-O-NEXT: Running pass: InlinerPass ; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass ; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass ; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass on (foo) diff --git a/llvm/test/Other/new-pm-print-pipeline.ll b/llvm/test/Other/new-pm-print-pipeline.ll index 6214d6abcef3a..c9b06ac7e6e8e 100644 --- a/llvm/test/Other/new-pm-print-pipeline.ll +++ b/llvm/test/Other/new-pm-print-pipeline.ll @@ -62,7 +62,7 @@ ; CHECK-20: cgscc(inline,inline),cgscc(inline) ; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='scc-oz-module-inliner' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-21 -; CHECK-21: require,function(invalidate),require,cgscc(devirt<4>(inline,inline,{{.*}},instcombine{{.*}})) +; CHECK-21: require,function(invalidate),require,cgscc(devirt<4>(inline,{{.*}},instcombine{{.*}})) ; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='cgscc(function(no-op-function)),function(no-op-function)' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-22 ; CHECK-22: cgscc(function(no-op-function)),function(no-op-function) diff --git a/llvm/test/Other/new-pm-thinlto-postlink-defaults.ll b/llvm/test/Other/new-pm-thinlto-postlink-defaults.ll index 1103c8f984fa2..064362eabbf83 100644 --- a/llvm/test/Other/new-pm-thinlto-postlink-defaults.ll +++ b/llvm/test/Other/new-pm-thinlto-postlink-defaults.ll @@ -61,6 +61,8 @@ ; CHECK-O-NEXT: Running analysis: TypeBasedAA ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy ; CHECK-O-NEXT: Running pass: SimplifyCFGPass +; CHECK-O-NEXT: Running pass: AlwaysInlinerPass +; CHECK-PRELINK-O-NEXT: Running analysis: ProfileSummaryAnalysis ; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass ; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA @@ -75,7 +77,6 @@ ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy ; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass ; CHECK-O-NEXT: Running pass: InlinerPass -; CHECK-O-NEXT: Running pass: InlinerPass ; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass ; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass ; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass on (foo) diff --git a/llvm/test/Other/new-pm-thinlto-postlink-pgo-defaults.ll b/llvm/test/Other/new-pm-thinlto-postlink-pgo-defaults.ll index 05d165ae399f9..19a44867e434a 100644 --- a/llvm/test/Other/new-pm-thinlto-postlink-pgo-defaults.ll +++ b/llvm/test/Other/new-pm-thinlto-postlink-pgo-defaults.ll @@ -50,6 +50,7 @@ ; CHECK-O-NEXT: Running analysis: LoopAnalysis on foo ; CHECK-O-NEXT: Running analysis: PostDominatorTreeAnalysis on foo ; CHECK-O-NEXT: Running pass: SimplifyCFGPass +; CHECK-O-NEXT: Running pass: AlwaysInlinerPass ; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass ; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA @@ -64,7 +65,6 @@ ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy<{{.*}}LazyCallGraph::SCC{{.*}}> ; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass ; CHECK-O-NEXT: Running pass: InlinerPass -; CHECK-O-NEXT: Running pass: InlinerPass ; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass ; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass ; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass diff --git a/llvm/test/Other/new-pm-thinlto-postlink-samplepgo-defaults.ll b/llvm/test/Other/new-pm-thinlto-postlink-samplepgo-defaults.ll index 2393ba9dbde4b..ac80a31d8fd4b 100644 --- a/llvm/test/Other/new-pm-thinlto-postlink-samplepgo-defaults.ll +++ b/llvm/test/Other/new-pm-thinlto-postlink-samplepgo-defaults.ll @@ -58,6 +58,8 @@ ; CHECK-O-NEXT: Running analysis: LoopAnalysis on foo ; CHECK-O-NEXT: Running analysis: PostDominatorTreeAnalysis on foo ; CHECK-O-NEXT: Running pass: SimplifyCFGPass on foo + +; CHECK-O-NEXT: Running pass: AlwaysInlinerPass ; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass ; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA @@ -71,7 +73,6 @@ ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy ; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass ; CHECK-O-NEXT: Running pass: InlinerPass -; CHECK-O-NEXT: Running pass: InlinerPass ; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass ; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass ; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass diff --git a/llvm/test/Other/new-pm-thinlto-prelink-defaults.ll b/llvm/test/Other/new-pm-thinlto-prelink-defaults.ll index 722d586c86907..6486639e07b49 100644 --- a/llvm/test/Other/new-pm-thinlto-prelink-defaults.ll +++ b/llvm/test/Other/new-pm-thinlto-prelink-defaults.ll @@ -92,6 +92,8 @@ ; CHECK-O-NEXT: Running analysis: TypeBasedAA ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy ; CHECK-O-NEXT: Running pass: SimplifyCFGPass +; CHECK-O-NEXT: Running pass: AlwaysInlinerPass +; CHECK-O-NEXT: Running analysis: ProfileSummaryAnalysis ; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass ; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA @@ -100,14 +102,12 @@ ; CHECK-O-NEXT: Running pass: InvalidateAnalysisPass<{{.*}}AAManager ; CHECK-O-NEXT: Invalidating analysis: AAManager ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}ProfileSummaryAnalysis -; CHECK-O-NEXT: Running analysis: ProfileSummaryAnalysis ; CHECK-O-NEXT: Running analysis: InnerAnalysisManagerProxy ; CHECK-O-NEXT: Running analysis: LazyCallGraphAnalysis ; CHECK-O-NEXT: Running analysis: FunctionAnalysisManagerCGSCCProxy ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy ; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass ; CHECK-O-NEXT: Running pass: InlinerPass -; CHECK-O-NEXT: Running pass: InlinerPass ; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass ; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass ; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass on (foo) diff --git a/llvm/test/Other/new-pm-thinlto-prelink-pgo-defaults.ll b/llvm/test/Other/new-pm-thinlto-prelink-pgo-defaults.ll index 7e4c82666f0cd..09f9f0f48badd 100644 --- a/llvm/test/Other/new-pm-thinlto-prelink-pgo-defaults.ll +++ b/llvm/test/Other/new-pm-thinlto-prelink-pgo-defaults.ll @@ -83,6 +83,7 @@ ; CHECK-O-NEXT: Running pass: PGOIndirectCallPromotion on ; CHECK-O-NEXT: Running analysis: InnerAnalysisManagerProxy ; CHECK-O-NEXT: Running analysis: OptimizationRemarkEmitterAnalysis on foo +; CHECK-O-NEXT: Running pass: AlwaysInlinerPass ; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass ; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA @@ -97,7 +98,6 @@ ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy<{{.*}}LazyCallGraph::SCC{{.*}}> ; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass ; CHECK-O-NEXT: Running pass: InlinerPass -; CHECK-O-NEXT: Running pass: InlinerPass ; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass ; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass ; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass diff --git a/llvm/test/Other/new-pm-thinlto-prelink-samplepgo-defaults.ll b/llvm/test/Other/new-pm-thinlto-prelink-samplepgo-defaults.ll index c609fbe31de21..47bdbfd2d357d 100644 --- a/llvm/test/Other/new-pm-thinlto-prelink-samplepgo-defaults.ll +++ b/llvm/test/Other/new-pm-thinlto-prelink-samplepgo-defaults.ll @@ -63,6 +63,7 @@ ; CHECK-O-NEXT: Running analysis: LoopAnalysis on foo ; CHECK-O-NEXT: Running analysis: PostDominatorTreeAnalysis on foo ; CHECK-O-NEXT: Running pass: SimplifyCFGPass on foo +; CHECK-O-NEXT: Running pass: AlwaysInlinerPass ; CHECK-O-NEXT: Running pass: ModuleInlinerWrapperPass ; CHECK-O-NEXT: Running analysis: InlineAdvisorAnalysis ; CHECK-O-NEXT: Running pass: RequireAnalysisPass<{{.*}}GlobalsAA @@ -76,7 +77,6 @@ ; CHECK-O-NEXT: Running analysis: OuterAnalysisManagerProxy<{{.*}}LazyCallGraph::SCC{{.*}}> ; CHECK-O-NEXT: Running pass: DevirtSCCRepeatedPass ; CHECK-O-NEXT: Running pass: InlinerPass -; CHECK-O-NEXT: Running pass: InlinerPass ; CHECK-O-NEXT: Running pass: PostOrderFunctionAttrsPass ; CHECK-O3-NEXT: Running pass: ArgumentPromotionPass ; CHECK-O2-NEXT: Running pass: OpenMPOptCGSCCPass diff --git a/llvm/test/Transforms/Inline/always-inline-phase-ordering.ll b/llvm/test/Transforms/Inline/always-inline-phase-ordering.ll new file mode 100644 index 0000000000000..e69ca48344900 --- /dev/null +++ b/llvm/test/Transforms/Inline/always-inline-phase-ordering.ll @@ -0,0 +1,164 @@ +; RUN: opt --Os -pass-remarks=inline -S < %s 2>&1 | FileCheck %s +target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128" +target triple = "arm64e-apple-macosx13" + +; CHECK: remark: :0:0: 'wibble' inlined into 'bar.8' with (cost=always): always inline attribute +; CHECK: remark: :0:0: 'wibble' inlined into 'pluto' with (cost=always): always inline attribute +; CHECK: remark: :0:0: 'snork' inlined into 'blam' with (cost=always): always inline attribute +; CHECK: remark: :0:0: 'wobble' inlined into 'blam' with (cost=always): always inline attribute +; CHECK: remark: :0:0: 'wobble' inlined into 'snork' with (cost=always): always inline attribute +; CHECK: remark: :0:0: 'spam' inlined into 'blam' with (cost=65, threshold=75) +; CHECK: remark: :0:0: 'wibble.1' inlined into 'widget' with (cost=30, threshold=75) +; CHECK: remark: :0:0: 'widget' inlined into 'bar.8' with (cost=30, threshold=75) +; CHECK: remark: :0:0: 'barney' inlined into 'wombat' with (cost=30, threshold=75) + +define linkonce_odr void @wombat(ptr %arg) #0 { +bb: + call void @barney() + ret void +} + +define i1 @foo() { +bb: + call void @wombat(ptr null) + unreachable +} + +define linkonce_odr void @pluto() #1 !prof !38 { +bb: + call void @wibble() + ret void +} + +; Function Attrs: alwaysinline +define linkonce_odr void @wibble() #2 { +bb: + call void @widget() + ret void +} + +define linkonce_odr void @widget() { +bb: + call void @wibble.1() + ret void +} + +define linkonce_odr void @wibble.1() { +bb: + %0 = call i32 @foo.2() + call void @blam() + ret void +} + +declare i32 @foo.2() + +define linkonce_odr void @blam() { +bb: + %tmp = call i32 @snork() + %tmpv1 = call ptr @wombat.3() + call void @eggs() + %tmpv2 = call ptr @wombat.3() + ret void +} + +; Function Attrs: alwaysinline +define linkonce_odr i32 @snork() #2 { +bb: + %tmpv1 = call i32 @spam() + %tmpv2 = call i32 @wobble() + call void @widget.4(i32 %tmpv2) + ret i32 0 +} + +declare void @eggs() + +declare ptr @wombat.3() + +define linkonce_odr i32 @spam() { +bb: + %tmpv1 = call i32 @wombat.6() + %tmpv2 = call i64 @wobble.5(i8 0) + %tmpv3 = call i64 @bar() + ret i32 0 +} + +; Function Attrs: alwaysinline +define linkonce_odr i32 @wobble() #2 { +bb: + %tmpv = call i64 @wobble.5(i8 0) + %tmpv1 = call i64 @eggs.7() + %tmpv2 = call i64 @wobble.5(i8 0) + %tmpv3 = call i64 @eggs.7() + %tmpv4 = lshr i64 %tmpv1, 1 + %tmpv5 = trunc i64 %tmpv4 to i32 + %tmpv6 = xor i32 %tmpv5, 23 + ret i32 %tmpv6 +} + +declare void @widget.4(i32) + +declare i64 @bar() + +declare i64 @wobble.5(i8) + +declare i32 @wombat.6() + +declare i64 @eggs.7() + +define linkonce_odr void @barney() { +bb: + call void @bar.8() + call void @pluto() + unreachable +} + +define linkonce_odr void @bar.8() { +bb: + call void @wibble() + ret void +} + +attributes #0 = { "frame-pointer"="non-leaf" } +attributes #1 = { "target-cpu"="apple-m1" } +attributes #2 = { alwaysinline } + +!llvm.module.flags = !{!0, !1, !30, !31, !32, !36, !37} + +!0 = !{i32 2, !"SDK Version", [2 x i32] [i32 13, i32 3]} +!1 = !{i32 1, !"ProfileSummary", !2} +!2 = !{!3, !4, !5, !6, !7, !8, !9, !10, !11, !12} +!3 = !{!"ProfileFormat", !"InstrProf"} +!4 = !{!"TotalCount", i64 864540306756} +!5 = !{!"MaxCount", i64 6596759955} +!6 = !{!"MaxInternalCount", i64 2828618424} +!7 = !{!"MaxFunctionCount", i64 6596759955} +!8 = !{!"NumCounts", i64 268920} +!9 = !{!"NumFunctions", i64 106162} +!10 = !{!"IsPartialProfile", i64 0} +!11 = !{!"PartialProfileRatio", double 0.000000e+00} +!12 = !{!"DetailedSummary", !13} +!13 = !{!14, !15, !16, !17, !18, !19, !20, !21, !22, !23, !24, !25, !26, !27, !28, !29} +!14 = !{i32 10000, i64 5109654023, i32 2} +!15 = !{i32 100000, i64 2480859832, i32 25} +!16 = !{i32 200000, i64 1566552109, i32 70} +!17 = !{i32 300000, i64 973667919, i32 140} +!18 = !{i32 400000, i64 552159773, i32 263} +!19 = !{i32 500000, i64 353879860, i32 463} +!20 = !{i32 600000, i64 187122455, i32 799} +!21 = !{i32 700000, i64 105465980, i32 1419} +!22 = !{i32 800000, i64 49243829, i32 2620} +!23 = !{i32 900000, i64 15198227, i32 5898} +!24 = !{i32 950000, i64 5545670, i32 10696} +!25 = !{i32 990000, i64 804816, i32 25738} +!26 = !{i32 999000, i64 73999, i32 53382} +!27 = !{i32 999900, i64 6530, i32 83503} +!28 = !{i32 999990, i64 899, i32 110416} +!29 = !{i32 999999, i64 120, i32 130201} +!30 = !{i32 7, !"Dwarf Version", i32 4} +!31 = !{i32 2, !"Debug Info Version", i32 3} +!32 = !{i32 1, !"wchar_size", i32 4} +!34 = !{!35} +!35 = !{i32 0, i1 false} +!36 = !{i32 8, !"PIC Level", i32 2} +!37 = !{i32 7, !"frame-pointer", i32 1} +!38 = !{!"function_entry_count", i64 15128150}