Skip to content

Commit c17f9c3

Browse files
nikicnikic
committed
[InstCombine] Support folding intrinsics into phis
1 parent 1a3e857 commit c17f9c3

File tree

5 files changed

+77
-66
lines changed

5 files changed

+77
-66
lines changed

‎clang/test/Headers/__clang_hip_math.hip

Lines changed: 60 additions & 40 deletions
Original file line numberDiff line numberDiff line change
@@ -4981,11 +4981,13 @@ extern "C" __device__ double test_normcdfinv(double x) {
49814981
// DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
49824982
// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 4
49834983
// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
4984-
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
4984+
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
4985+
// DEFAULT: _ZL5normfiPKf.exit.loopexit:
4986+
// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[ADD_I]])
4987+
// DEFAULT-NEXT: br label [[_ZL5NORMFIPKF_EXIT]]
49854988
// DEFAULT: _ZL5normfiPKf.exit:
4986-
// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
4987-
// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[__R_0_I_LCSSA]])
4988-
// DEFAULT-NEXT: ret float [[TMP1]]
4989+
// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL5NORMFIPKF_EXIT_LOOPEXIT]] ]
4990+
// DEFAULT-NEXT: ret float [[__R_0_I_LCSSA]]
49894991
//
49904992
// FINITEONLY-LABEL: @test_normf(
49914993
// FINITEONLY-NEXT: entry:
@@ -5001,11 +5003,13 @@ extern "C" __device__ double test_normcdfinv(double x) {
50015003
// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_0_I4]], [[MUL_I]]
50025004
// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 4
50035005
// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5004-
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
5006+
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
5007+
// FINITEONLY: _ZL5normfiPKf.exit.loopexit:
5008+
// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract noundef float @llvm.sqrt.f32(float [[ADD_I]])
5009+
// FINITEONLY-NEXT: br label [[_ZL5NORMFIPKF_EXIT]]
50055010
// FINITEONLY: _ZL5normfiPKf.exit:
5006-
// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
5007-
// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract noundef float @llvm.sqrt.f32(float [[__R_0_I_LCSSA]])
5008-
// FINITEONLY-NEXT: ret float [[TMP1]]
5011+
// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL5NORMFIPKF_EXIT_LOOPEXIT]] ]
5012+
// FINITEONLY-NEXT: ret float [[__R_0_I_LCSSA]]
50095013
//
50105014
// APPROX-LABEL: @test_normf(
50115015
// APPROX-NEXT: entry:
@@ -5021,11 +5025,13 @@ extern "C" __device__ double test_normcdfinv(double x) {
50215025
// APPROX-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
50225026
// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 4
50235027
// APPROX-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5024-
// APPROX-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
5028+
// APPROX-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
5029+
// APPROX: _ZL5normfiPKf.exit.loopexit:
5030+
// APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[ADD_I]])
5031+
// APPROX-NEXT: br label [[_ZL5NORMFIPKF_EXIT]]
50255032
// APPROX: _ZL5normfiPKf.exit:
5026-
// APPROX-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
5027-
// APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[__R_0_I_LCSSA]])
5028-
// APPROX-NEXT: ret float [[TMP1]]
5033+
// APPROX-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL5NORMFIPKF_EXIT_LOOPEXIT]] ]
5034+
// APPROX-NEXT: ret float [[__R_0_I_LCSSA]]
50295035
//
50305036
// NCRDIV-LABEL: @test_normf(
50315037
// NCRDIV-NEXT: entry:
@@ -5041,11 +5047,13 @@ extern "C" __device__ double test_normcdfinv(double x) {
50415047
// NCRDIV-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
50425048
// NCRDIV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 4
50435049
// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5044-
// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
5050+
// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
5051+
// NCRDIV: _ZL5normfiPKf.exit.loopexit:
5052+
// NCRDIV-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[ADD_I]]), !fpmath [[META22:![0-9]+]]
5053+
// NCRDIV-NEXT: br label [[_ZL5NORMFIPKF_EXIT]]
50455054
// NCRDIV: _ZL5normfiPKf.exit:
5046-
// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
5047-
// NCRDIV-NEXT: [[TMP1:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[__R_0_I_LCSSA]]), !fpmath [[META22:![0-9]+]]
5048-
// NCRDIV-NEXT: ret float [[TMP1]]
5055+
// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL5NORMFIPKF_EXIT_LOOPEXIT]] ]
5056+
// NCRDIV-NEXT: ret float [[__R_0_I_LCSSA]]
50495057
//
50505058
// AMDGCNSPIRV-LABEL: @test_normf(
50515059
// AMDGCNSPIRV-NEXT: entry:
@@ -5061,11 +5069,13 @@ extern "C" __device__ double test_normcdfinv(double x) {
50615069
// AMDGCNSPIRV-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
50625070
// AMDGCNSPIRV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__A_ADDR_0_I3]], i64 4
50635071
// AMDGCNSPIRV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5064-
// AMDGCNSPIRV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
5072+
// AMDGCNSPIRV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
5073+
// AMDGCNSPIRV: _ZL5normfiPKf.exit.loopexit:
5074+
// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = tail call contract noundef addrspace(4) float @llvm.sqrt.f32(float [[ADD_I]])
5075+
// AMDGCNSPIRV-NEXT: br label [[_ZL5NORMFIPKF_EXIT]]
50655076
// AMDGCNSPIRV: _ZL5normfiPKf.exit:
5066-
// AMDGCNSPIRV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
5067-
// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = tail call contract noundef addrspace(4) float @llvm.sqrt.f32(float [[__R_0_I_LCSSA]])
5068-
// AMDGCNSPIRV-NEXT: ret float [[TMP1]]
5077+
// AMDGCNSPIRV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL5NORMFIPKF_EXIT_LOOPEXIT]] ]
5078+
// AMDGCNSPIRV-NEXT: ret float [[__R_0_I_LCSSA]]
50695079
//
50705080
extern "C" __device__ float test_normf(int x, const float *y) {
50715081
return normf(x, y);
@@ -5085,11 +5095,13 @@ extern "C" __device__ float test_normf(int x, const float *y) {
50855095
// DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
50865096
// DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 8
50875097
// DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5088-
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
5098+
// DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
5099+
// DEFAULT: _ZL4normiPKd.exit.loopexit:
5100+
// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[ADD_I]])
5101+
// DEFAULT-NEXT: br label [[_ZL4NORMIPKD_EXIT]]
50895102
// DEFAULT: _ZL4normiPKd.exit:
5090-
// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
5091-
// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[__R_0_I_LCSSA]])
5092-
// DEFAULT-NEXT: ret double [[TMP1]]
5103+
// DEFAULT-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL4NORMIPKD_EXIT_LOOPEXIT]] ]
5104+
// DEFAULT-NEXT: ret double [[__R_0_I_LCSSA]]
50935105
//
50945106
// FINITEONLY-LABEL: @test_norm(
50955107
// FINITEONLY-NEXT: entry:
@@ -5105,11 +5117,13 @@ extern "C" __device__ float test_normf(int x, const float *y) {
51055117
// FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_0_I4]], [[MUL_I]]
51065118
// FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 8
51075119
// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5108-
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
5120+
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
5121+
// FINITEONLY: _ZL4normiPKd.exit.loopexit:
5122+
// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract noundef double @llvm.sqrt.f64(double [[ADD_I]])
5123+
// FINITEONLY-NEXT: br label [[_ZL4NORMIPKD_EXIT]]
51095124
// FINITEONLY: _ZL4normiPKd.exit:
5110-
// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
5111-
// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract noundef double @llvm.sqrt.f64(double [[__R_0_I_LCSSA]])
5112-
// FINITEONLY-NEXT: ret double [[TMP1]]
5125+
// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL4NORMIPKD_EXIT_LOOPEXIT]] ]
5126+
// FINITEONLY-NEXT: ret double [[__R_0_I_LCSSA]]
51135127
//
51145128
// APPROX-LABEL: @test_norm(
51155129
// APPROX-NEXT: entry:
@@ -5125,11 +5139,13 @@ extern "C" __device__ float test_normf(int x, const float *y) {
51255139
// APPROX-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
51265140
// APPROX-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 8
51275141
// APPROX-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5128-
// APPROX-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
5142+
// APPROX-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
5143+
// APPROX: _ZL4normiPKd.exit.loopexit:
5144+
// APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[ADD_I]])
5145+
// APPROX-NEXT: br label [[_ZL4NORMIPKD_EXIT]]
51295146
// APPROX: _ZL4normiPKd.exit:
5130-
// APPROX-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
5131-
// APPROX-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[__R_0_I_LCSSA]])
5132-
// APPROX-NEXT: ret double [[TMP1]]
5147+
// APPROX-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL4NORMIPKD_EXIT_LOOPEXIT]] ]
5148+
// APPROX-NEXT: ret double [[__R_0_I_LCSSA]]
51335149
//
51345150
// NCRDIV-LABEL: @test_norm(
51355151
// NCRDIV-NEXT: entry:
@@ -5145,11 +5161,13 @@ extern "C" __device__ float test_normf(int x, const float *y) {
51455161
// NCRDIV-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
51465162
// NCRDIV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 8
51475163
// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5148-
// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
5164+
// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
5165+
// NCRDIV: _ZL4normiPKd.exit.loopexit:
5166+
// NCRDIV-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[ADD_I]])
5167+
// NCRDIV-NEXT: br label [[_ZL4NORMIPKD_EXIT]]
51495168
// NCRDIV: _ZL4normiPKd.exit:
5150-
// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
5151-
// NCRDIV-NEXT: [[TMP1:%.*]] = tail call contract noundef double @llvm.sqrt.f64(double [[__R_0_I_LCSSA]])
5152-
// NCRDIV-NEXT: ret double [[TMP1]]
5169+
// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL4NORMIPKD_EXIT_LOOPEXIT]] ]
5170+
// NCRDIV-NEXT: ret double [[__R_0_I_LCSSA]]
51535171
//
51545172
// AMDGCNSPIRV-LABEL: @test_norm(
51555173
// AMDGCNSPIRV-NEXT: entry:
@@ -5165,11 +5183,13 @@ extern "C" __device__ float test_normf(int x, const float *y) {
51655183
// AMDGCNSPIRV-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
51665184
// AMDGCNSPIRV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr addrspace(4) [[__A_ADDR_0_I3]], i64 8
51675185
// AMDGCNSPIRV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5168-
// AMDGCNSPIRV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
5186+
// AMDGCNSPIRV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT_LOOPEXIT:%.*]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
5187+
// AMDGCNSPIRV: _ZL4normiPKd.exit.loopexit:
5188+
// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = tail call contract noundef addrspace(4) double @llvm.sqrt.f64(double [[ADD_I]])
5189+
// AMDGCNSPIRV-NEXT: br label [[_ZL4NORMIPKD_EXIT]]
51695190
// AMDGCNSPIRV: _ZL4normiPKd.exit:
5170-
// AMDGCNSPIRV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
5171-
// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = tail call contract noundef addrspace(4) double @llvm.sqrt.f64(double [[__R_0_I_LCSSA]])
5172-
// AMDGCNSPIRV-NEXT: ret double [[TMP1]]
5191+
// AMDGCNSPIRV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[TMP1]], [[_ZL4NORMIPKD_EXIT_LOOPEXIT]] ]
5192+
// AMDGCNSPIRV-NEXT: ret double [[__R_0_I_LCSSA]]
51735193
//
51745194
extern "C" __device__ double test_norm(int x, const double *y) {
51755195
return norm(x, y);

‎llvm/lib/Transforms/InstCombine/InstCombineCalls.cpp

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3897,10 +3897,14 @@ Instruction *InstCombinerImpl::visitCallInst(CallInst &CI) {
38973897
// perform cross-lane operations.
38983898
if (isSafeToSpeculativelyExecuteWithVariableReplaced(&CI) &&
38993899
isNotCrossLaneOperation(II))
3900-
for (Value *Op : II->args())
3900+
for (Value *Op : II->args()) {
39013901
if (auto *Sel = dyn_cast<SelectInst>(Op))
39023902
if (Instruction *R = FoldOpIntoSelect(*II, Sel))
39033903
return R;
3904+
if (auto *Phi = dyn_cast<PHINode>(Op))
3905+
if (Instruction *R = foldOpIntoPhi(*II, Phi))
3906+
return R;
3907+
}
39043908

39053909
if (Instruction *Shuf = foldShuffledIntrinsicOperands(II))
39063910
return Shuf;

‎llvm/test/Transforms/InstCombine/fpclass-from-dom-cond.ll

Lines changed: 3 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -131,10 +131,10 @@ define i1 @test5(double %x, i1 %cond) {
131131
; CHECK: if.then:
132132
; CHECK-NEXT: ret i1 false
133133
; CHECK: if.end:
134+
; CHECK-NEXT: [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f64(double [[X]], i32 408)
134135
; CHECK-NEXT: br label [[EXIT]]
135136
; CHECK: exit:
136-
; CHECK-NEXT: [[Y:%.*]] = phi double [ -1.000000e+00, [[ENTRY:%.*]] ], [ [[X]], [[IF_END]] ]
137-
; CHECK-NEXT: [[RET:%.*]] = tail call i1 @llvm.is.fpclass.f64(double [[Y]], i32 408)
137+
; CHECK-NEXT: [[RET:%.*]] = phi i1 [ true, [[ENTRY:%.*]] ], [ [[TMP0]], [[IF_END]] ]
138138
; CHECK-NEXT: ret i1 [[RET]]
139139
;
140140
entry:
@@ -391,11 +391,9 @@ define float @test_signbit_check_fail(float %x, i1 %cond) {
391391
; CHECK: if.else:
392392
; CHECK-NEXT: br i1 [[COND]], label [[IF_THEN2:%.*]], label [[IF_END]]
393393
; CHECK: if.then2:
394-
; CHECK-NEXT: [[FNEG2:%.*]] = fneg float [[X]]
395394
; CHECK-NEXT: br label [[IF_END]]
396395
; CHECK: if.end:
397-
; CHECK-NEXT: [[VALUE:%.*]] = phi float [ [[FNEG]], [[IF_THEN1]] ], [ [[FNEG2]], [[IF_THEN2]] ], [ [[X]], [[IF_ELSE]] ]
398-
; CHECK-NEXT: [[RET:%.*]] = call float @llvm.fabs.f32(float [[VALUE]])
396+
; CHECK-NEXT: [[RET:%.*]] = phi float [ [[FNEG]], [[IF_THEN1]] ], [ [[X]], [[IF_THEN2]] ], [ [[X]], [[IF_ELSE]] ]
399397
; CHECK-NEXT: ret float [[RET]]
400398
;
401399
%i32 = bitcast float %x to i32

‎llvm/test/Transforms/InstCombine/known-phi-recurse.ll

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -261,14 +261,11 @@ define i8 @knownbits_umax_select_test() {
261261
; CHECK-NEXT: entry:
262262
; CHECK-NEXT: br label [[LOOP:%.*]]
263263
; CHECK: loop:
264-
; CHECK-NEXT: [[INDVAR:%.*]] = phi i8 [ 0, [[ENTRY:%.*]] ], [ [[CONTAIN:%.*]], [[LOOP]] ]
265264
; CHECK-NEXT: [[COND0:%.*]] = call i1 @cond()
266-
; CHECK-NEXT: [[CONTAIN]] = call i8 @llvm.umax.i8(i8 [[INDVAR]], i8 1)
267265
; CHECK-NEXT: [[COND1:%.*]] = call i1 @cond()
268266
; CHECK-NEXT: br i1 [[COND1]], label [[EXIT:%.*]], label [[LOOP]]
269267
; CHECK: exit:
270-
; CHECK-NEXT: [[BOOL:%.*]] = and i8 [[CONTAIN]], 1
271-
; CHECK-NEXT: ret i8 [[BOOL]]
268+
; CHECK-NEXT: ret i8 1
272269
;
273270
entry:
274271
br label %loop

0 commit comments

Comments
 (0)