Skip to content

Commit 16d7383

Browse files
nikicnikic
authored
[InstCombine] Support folding intrinsics into phis (#151115)
Call foldOpIntoPhi() for speculatable intrinsics. We already do this for FoldOpIntoSelect(). Among other things, this partially subsumes #149858.
1 parent 9a9b8b7 commit 16d7383

File tree

7 files changed

+110
-70
lines changed

7 files changed

+110
-70
lines changed

‎clang/test/Headers/__clang_hip_math.hip‎

Lines changed: 63 additions & 43 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 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 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 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 float @llvm.sqrt.f32(float [[ADD_I]])
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 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 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 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 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 [[LOOP22:![0-9]+]]
5165+
// NCRDIV: _ZL4normiPKd.exit.loopexit:
5166+
// NCRDIV-NEXT: [[TMP1:%.*]] = tail call contract 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 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);
@@ -5707,7 +5727,7 @@ extern "C" __device__ double test_rint(double x) {
57075727
// NCRDIV-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
57085728
// NCRDIV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 4
57095729
// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5710-
// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
5730+
// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
57115731
// NCRDIV: _ZL6rnormfiPKf.exit:
57125732
// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
57135733
// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rsqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR15]]
@@ -5811,7 +5831,7 @@ extern "C" __device__ float test_rnormf(int x, const float* y) {
58115831
// NCRDIV-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
58125832
// NCRDIV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 8
58135833
// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
5814-
// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]]
5834+
// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
58155835
// NCRDIV: _ZL5rnormiPKd.exit:
58165836
// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
58175837
// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rsqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR15]]
@@ -6616,7 +6636,7 @@ extern "C" __device__ double test_sinpi(double x) {
66166636
//
66176637
// NCRDIV-LABEL: @test_sqrtf(
66186638
// NCRDIV-NEXT: entry:
6619-
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[X:%.*]]), !fpmath [[META22]]
6639+
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[X:%.*]]), !fpmath [[META25:![0-9]+]]
66206640
// NCRDIV-NEXT: ret float [[TMP0]]
66216641
//
66226642
// AMDGCNSPIRV-LABEL: @test_sqrtf(

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

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3891,16 +3891,20 @@ Instruction *InstCombinerImpl::visitCallInst(CallInst &CI) {
38913891
}
38923892
}
38933893

3894-
// Try to fold intrinsic into select operands. This is legal if:
3894+
// Try to fold intrinsic into select/phi operands. This is legal if:
38953895
// * The intrinsic is speculatable.
38963896
// * The select condition is not a vector, or the intrinsic does not
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/lib/Transforms/InstCombine/InstructionCombining.cpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1994,6 +1994,8 @@ Instruction *InstCombinerImpl::foldOpIntoPhi(Instruction &I, PHINode *PN,
19941994
}
19951995
Clone = InsertNewInstBefore(Clone, OpBB->getTerminator()->getIterator());
19961996
Clones.insert({OpBB, Clone});
1997+
// We may have speculated the instruction.
1998+
Clone->dropUBImplyingAttrsAndMetadata();
19971999
}
19982000

19992001
NewPhiValues[OpIndex] = Clone;

‎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

‎llvm/test/Transforms/InstCombine/phi.ll‎

Lines changed: 27 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2998,3 +2998,30 @@ join:
29982998
%cmp = icmp eq i32 %phi, 0
29992999
ret i1 %cmp
30003000
}
3001+
3002+
declare void @may_exit()
3003+
3004+
define i32 @intrinsic_over_phi_noundef(i1 %c, i1 %c2, i32 %a) {
3005+
; CHECK-LABEL: @intrinsic_over_phi_noundef(
3006+
; CHECK-NEXT: entry:
3007+
; CHECK-NEXT: br i1 [[C:%.*]], label [[IF:%.*]], label [[JOIN:%.*]]
3008+
; CHECK: if:
3009+
; CHECK-NEXT: [[TMP0:%.*]] = call i32 @llvm.umax.i32(i32 [[A:%.*]], i32 1)
3010+
; CHECK-NEXT: br label [[JOIN]]
3011+
; CHECK: join:
3012+
; CHECK-NEXT: [[PHI:%.*]] = phi i32 [ [[TMP0]], [[IF]] ], [ 1, [[ENTRY:%.*]] ]
3013+
; CHECK-NEXT: call void @may_exit()
3014+
; CHECK-NEXT: ret i32 [[PHI]]
3015+
;
3016+
entry:
3017+
br i1 %c, label %if, label %join
3018+
3019+
if:
3020+
br label %join
3021+
3022+
join:
3023+
%phi = phi i32 [ %a, %if ], [ 0, %entry ]
3024+
call void @may_exit()
3025+
%umax = call noundef i32 @llvm.umax(i32 noundef %phi, i32 1)
3026+
ret i32 %umax
3027+
}

0 commit comments

Comments
 (0)