Skip to content

Commit b6bd41d

Browse files
goldsteinngoldsteinn
committed
[InstCombine] Add canonicalization of sitofp -> uitofp nneg
This is essentially the same as #82404 but has the `nneg` flag which allows the backend to reliably undo the transform. Closes #88299
1 parent 34777c2 commit b6bd41d

File tree

12 files changed

+127
-113
lines changed

12 files changed

+127
-113
lines changed

‎clang/test/Headers/__clang_hip_math.hip‎

Lines changed: 12 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -1685,7 +1685,7 @@ extern "C" __device__ double test_j1(double x) {
16851685
// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
16861686
// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
16871687
// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
1688-
// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
1688+
// DEFAULT-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to float
16891689
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
16901690
// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
16911691
// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
@@ -1718,7 +1718,7 @@ extern "C" __device__ double test_j1(double x) {
17181718
// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
17191719
// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
17201720
// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
1721-
// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
1721+
// FINITEONLY-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to float
17221722
// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]]
17231723
// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_0_I3]], [[DIV_I]]
17241724
// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_0_I2]]
@@ -1751,7 +1751,7 @@ extern "C" __device__ double test_j1(double x) {
17511751
// APPROX-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
17521752
// APPROX-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
17531753
// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
1754-
// APPROX-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
1754+
// APPROX-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to float
17551755
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
17561756
// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
17571757
// APPROX-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
@@ -1788,7 +1788,7 @@ extern "C" __device__ float test_jnf(int x, float y) {
17881788
// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
17891789
// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
17901790
// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
1791-
// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
1791+
// DEFAULT-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to double
17921792
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
17931793
// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
17941794
// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]
@@ -1821,7 +1821,7 @@ extern "C" __device__ float test_jnf(int x, float y) {
18211821
// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
18221822
// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
18231823
// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
1824-
// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
1824+
// FINITEONLY-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to double
18251825
// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]]
18261826
// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_0_I3]], [[DIV_I]]
18271827
// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_0_I2]]
@@ -1854,7 +1854,7 @@ extern "C" __device__ float test_jnf(int x, float y) {
18541854
// APPROX-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
18551855
// APPROX-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
18561856
// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
1857-
// APPROX-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
1857+
// APPROX-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to double
18581858
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
18591859
// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
18601860
// APPROX-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]
@@ -4222,7 +4222,7 @@ extern "C" __device__ double test_y1(double x) {
42224222
// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
42234223
// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
42244224
// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
4225-
// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
4225+
// DEFAULT-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to float
42264226
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
42274227
// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
42284228
// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
@@ -4255,7 +4255,7 @@ extern "C" __device__ double test_y1(double x) {
42554255
// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
42564256
// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
42574257
// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
4258-
// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
4258+
// FINITEONLY-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to float
42594259
// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]]
42604260
// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_0_I3]], [[DIV_I]]
42614261
// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_0_I2]]
@@ -4288,7 +4288,7 @@ extern "C" __device__ double test_y1(double x) {
42884288
// APPROX-NEXT: [[__X1_0_I3:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
42894289
// APPROX-NEXT: [[__X0_0_I2:%.*]] = phi float [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
42904290
// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
4291-
// APPROX-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
4291+
// APPROX-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to float
42924292
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
42934293
// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
42944294
// APPROX-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
@@ -4325,7 +4325,7 @@ extern "C" __device__ float test_ynf(int x, float y) {
43254325
// DEFAULT-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
43264326
// DEFAULT-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
43274327
// DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
4328-
// DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
4328+
// DEFAULT-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to double
43294329
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
43304330
// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
43314331
// DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]
@@ -4358,7 +4358,7 @@ extern "C" __device__ float test_ynf(int x, float y) {
43584358
// FINITEONLY-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
43594359
// FINITEONLY-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
43604360
// FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
4361-
// FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
4361+
// FINITEONLY-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to double
43624362
// FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]]
43634363
// FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_0_I3]], [[DIV_I]]
43644364
// FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_0_I2]]
@@ -4391,7 +4391,7 @@ extern "C" __device__ float test_ynf(int x, float y) {
43914391
// APPROX-NEXT: [[__X1_0_I3:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
43924392
// APPROX-NEXT: [[__X0_0_I2:%.*]] = phi double [ [[__X1_0_I3]], [[FOR_BODY_I]] ], [ [[CALL_I_I]], [[IF_END4_I]] ]
43934393
// APPROX-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_0_I4]], 1
4394-
// APPROX-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
4394+
// APPROX-NEXT: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to double
43954395
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
43964396
// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
43974397
// APPROX-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]

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

Lines changed: 16 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1977,11 +1977,25 @@ Instruction *InstCombinerImpl::visitFPToSI(FPToSIInst &FI) {
19771977
}
19781978

19791979
Instruction *InstCombinerImpl::visitUIToFP(CastInst &CI) {
1980-
return commonCastTransforms(CI);
1980+
if (Instruction *R = commonCastTransforms(CI))
1981+
return R;
1982+
if (!CI.hasNonNeg() && isKnownNonNegative(CI.getOperand(0), SQ)) {
1983+
CI.setNonNeg();
1984+
return &CI;
1985+
}
1986+
return nullptr;
19811987
}
19821988

19831989
Instruction *InstCombinerImpl::visitSIToFP(CastInst &CI) {
1984-
return commonCastTransforms(CI);
1990+
if (Instruction *R = commonCastTransforms(CI))
1991+
return R;
1992+
if (isKnownNonNegative(CI.getOperand(0), SQ)) {
1993+
auto UI =
1994+
CastInst::Create(Instruction::UIToFP, CI.getOperand(0), CI.getType());
1995+
UI->setNonNeg(true);
1996+
return UI;
1997+
}
1998+
return nullptr;
19851999
}
19862000

19872001
Instruction *InstCombinerImpl::visitIntToPtr(IntToPtrInst &CI) {

‎llvm/test/Transforms/InstCombine/add-sitofp.ll‎

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -6,7 +6,7 @@ define double @x(i32 %a, i32 %b) {
66
; CHECK-NEXT: [[M:%.*]] = lshr i32 [[A:%.*]], 24
77
; CHECK-NEXT: [[N:%.*]] = and i32 [[M]], [[B:%.*]]
88
; CHECK-NEXT: [[TMP1:%.*]] = add nuw nsw i32 [[N]], 1
9-
; CHECK-NEXT: [[P:%.*]] = uitofp i32 [[TMP1]] to double
9+
; CHECK-NEXT: [[P:%.*]] = uitofp nneg i32 [[TMP1]] to double
1010
; CHECK-NEXT: ret double [[P]]
1111
;
1212
%m = lshr i32 %a, 24
@@ -20,7 +20,7 @@ define double @test(i32 %a) {
2020
; CHECK-LABEL: @test(
2121
; CHECK-NEXT: [[A_AND:%.*]] = and i32 [[A:%.*]], 1073741823
2222
; CHECK-NEXT: [[TMP1:%.*]] = add nuw nsw i32 [[A_AND]], 1
23-
; CHECK-NEXT: [[RES:%.*]] = uitofp i32 [[TMP1]] to double
23+
; CHECK-NEXT: [[RES:%.*]] = uitofp nneg i32 [[TMP1]] to double
2424
; CHECK-NEXT: ret double [[RES]]
2525
;
2626
; Drop two highest bits to guarantee that %a + 1 doesn't overflow
@@ -33,7 +33,7 @@ define double @test(i32 %a) {
3333
define float @test_neg(i32 %a) {
3434
; CHECK-LABEL: @test_neg(
3535
; CHECK-NEXT: [[A_AND:%.*]] = and i32 [[A:%.*]], 1073741823
36-
; CHECK-NEXT: [[A_AND_FP:%.*]] = sitofp i32 [[A_AND]] to float
36+
; CHECK-NEXT: [[A_AND_FP:%.*]] = uitofp nneg i32 [[A_AND]] to float
3737
; CHECK-NEXT: [[RES:%.*]] = fadd float [[A_AND_FP]], 1.000000e+00
3838
; CHECK-NEXT: ret float [[RES]]
3939
;
@@ -49,7 +49,7 @@ define double @test_2(i32 %a, i32 %b) {
4949
; CHECK-NEXT: [[A_AND:%.*]] = and i32 [[A:%.*]], 1073741823
5050
; CHECK-NEXT: [[B_AND:%.*]] = and i32 [[B:%.*]], 1073741823
5151
; CHECK-NEXT: [[TMP1:%.*]] = add nuw nsw i32 [[A_AND]], [[B_AND]]
52-
; CHECK-NEXT: [[RES:%.*]] = uitofp i32 [[TMP1]] to double
52+
; CHECK-NEXT: [[RES:%.*]] = uitofp nneg i32 [[TMP1]] to double
5353
; CHECK-NEXT: ret double [[RES]]
5454
;
5555
; Drop two highest bits to guarantee that %a + %b doesn't overflow
@@ -67,8 +67,8 @@ define float @test_2_neg(i32 %a, i32 %b) {
6767
; CHECK-LABEL: @test_2_neg(
6868
; CHECK-NEXT: [[A_AND:%.*]] = and i32 [[A:%.*]], 1073741823
6969
; CHECK-NEXT: [[B_AND:%.*]] = and i32 [[B:%.*]], 1073741823
70-
; CHECK-NEXT: [[A_AND_FP:%.*]] = sitofp i32 [[A_AND]] to float
71-
; CHECK-NEXT: [[B_AND_FP:%.*]] = sitofp i32 [[B_AND]] to float
70+
; CHECK-NEXT: [[A_AND_FP:%.*]] = uitofp nneg i32 [[A_AND]] to float
71+
; CHECK-NEXT: [[B_AND_FP:%.*]] = uitofp nneg i32 [[B_AND]] to float
7272
; CHECK-NEXT: [[RES:%.*]] = fadd float [[A_AND_FP]], [[B_AND_FP]]
7373
; CHECK-NEXT: ret float [[RES]]
7474
;
@@ -89,7 +89,7 @@ define float @test_3(i32 %a, i32 %b) {
8989
; CHECK-NEXT: [[M:%.*]] = lshr i32 [[A:%.*]], 24
9090
; CHECK-NEXT: [[N:%.*]] = and i32 [[M]], [[B:%.*]]
9191
; CHECK-NEXT: [[TMP1:%.*]] = add nuw nsw i32 [[N]], 1
92-
; CHECK-NEXT: [[P:%.*]] = uitofp i32 [[TMP1]] to float
92+
; CHECK-NEXT: [[P:%.*]] = uitofp nneg i32 [[TMP1]] to float
9393
; CHECK-NEXT: ret float [[P]]
9494
;
9595
%m = lshr i32 %a, 24
@@ -104,7 +104,7 @@ define <4 x double> @test_4(<4 x i32> %a, <4 x i32> %b) {
104104
; CHECK-NEXT: [[A_AND:%.*]] = and <4 x i32> [[A:%.*]], <i32 1073741823, i32 1073741823, i32 1073741823, i32 1073741823>
105105
; CHECK-NEXT: [[B_AND:%.*]] = and <4 x i32> [[B:%.*]], <i32 1073741823, i32 1073741823, i32 1073741823, i32 1073741823>
106106
; CHECK-NEXT: [[TMP1:%.*]] = add nuw nsw <4 x i32> [[A_AND]], [[B_AND]]
107-
; CHECK-NEXT: [[RES:%.*]] = uitofp <4 x i32> [[TMP1]] to <4 x double>
107+
; CHECK-NEXT: [[RES:%.*]] = uitofp nneg <4 x i32> [[TMP1]] to <4 x double>
108108
; CHECK-NEXT: ret <4 x double> [[RES]]
109109
;
110110
; Drop two highest bits to guarantee that %a + %b doesn't overflow
@@ -122,8 +122,8 @@ define <4 x float> @test_4_neg(<4 x i32> %a, <4 x i32> %b) {
122122
; CHECK-LABEL: @test_4_neg(
123123
; CHECK-NEXT: [[A_AND:%.*]] = and <4 x i32> [[A:%.*]], <i32 1073741823, i32 1073741823, i32 1073741823, i32 1073741823>
124124
; CHECK-NEXT: [[B_AND:%.*]] = and <4 x i32> [[B:%.*]], <i32 1073741823, i32 1073741823, i32 1073741823, i32 1073741823>
125-
; CHECK-NEXT: [[A_AND_FP:%.*]] = sitofp <4 x i32> [[A_AND]] to <4 x float>
126-
; CHECK-NEXT: [[B_AND_FP:%.*]] = sitofp <4 x i32> [[B_AND]] to <4 x float>
125+
; CHECK-NEXT: [[A_AND_FP:%.*]] = uitofp nneg <4 x i32> [[A_AND]] to <4 x float>
126+
; CHECK-NEXT: [[B_AND_FP:%.*]] = uitofp nneg <4 x i32> [[B_AND]] to <4 x float>
127127
; CHECK-NEXT: [[RES:%.*]] = fadd <4 x float> [[A_AND_FP]], [[B_AND_FP]]
128128
; CHECK-NEXT: ret <4 x float> [[RES]]
129129
;

0 commit comments

Comments
 (0)