[CVP][SCCP] Add support for uitofp nneg
#86154
Closed
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
llvmbot
added
clang
|
@llvm/pr-subscribers-function-specialization @llvm/pr-subscribers-clang Author: None (goldsteinn) Changes
Patch is 83.64 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/86154.diff 18 Files Affected:
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index 37099de74fb8ec..bff1708120604b 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -1685,7 +1685,7 @@ extern "C" __device__ double test_j1(double x) {
// 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: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to float
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
@@ -1718,7 +1718,7 @@ extern "C" __device__ double test_j1(double x) {
// 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: [[CONV_I:%.*]] = uitofp nneg 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_0_I3]], [[DIV_I]]
// 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) {
// 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: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to float
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
// 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) {
// 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: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to double
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
// 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) {
// 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: [[CONV_I:%.*]] = uitofp nneg 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_0_I3]], [[DIV_I]]
// 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) {
// 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: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to double
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
// APPROX-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]
@@ -4222,7 +4222,7 @@ extern "C" __device__ double test_y1(double x) {
// 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: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to float
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
// DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_0_I2]]
@@ -4255,7 +4255,7 @@ extern "C" __device__ double test_y1(double x) {
// 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: [[CONV_I:%.*]] = uitofp nneg 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_0_I3]], [[DIV_I]]
// 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) {
// 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: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to float
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_0_I3]], [[DIV_I]]
// 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) {
// 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: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to double
// DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
// DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
// 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) {
// 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: [[CONV_I:%.*]] = uitofp nneg 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_0_I3]], [[DIV_I]]
// 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) {
// 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: [[CONV_I:%.*]] = uitofp nneg i32 [[MUL_I]] to double
// APPROX-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
// APPROX-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_0_I3]], [[DIV_I]]
// APPROX-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_0_I2]]
diff --git a/llvm/lib/Transforms/InstCombine/InstCombineCasts.cpp b/llvm/lib/Transforms/InstCombine/InstCombineCasts.cpp
index 089a70c6e6cca1..d2ee83ace7f124 100644
--- a/llvm/lib/Transforms/InstCombine/InstCombineCasts.cpp
+++ b/llvm/lib/Transforms/InstCombine/InstCombineCasts.cpp
@@ -1957,11 +1957,25 @@ Instruction *InstCombinerImpl::visitFPToSI(FPToSIInst &FI) {
}
Instruction *InstCombinerImpl::visitUIToFP(CastInst &CI) {
- return commonCastTransforms(CI);
+ if (Instruction *R = commonCastTransforms(CI))
+ return R;
+ if (!CI.hasNonNeg() && isKnownNonNegative(CI.getOperand(0), SQ)) {
+ CI.setNonNeg();
+ return &CI;
+ }
+ return nullptr;
}
Instruction *InstCombinerImpl::visitSIToFP(CastInst &CI) {
- return commonCastTransforms(CI);
+ if (Instruction *R = commonCastTransforms(CI))
+ return R;
+ if (isKnownNonNegative(CI.getOperand(0), SQ)) {
+ auto UI =
+ CastInst::Create(Instruction::UIToFP, CI.getOperand(0), CI.getType());
+ UI->setNonNeg(true);
+ return UI;
+ }
+ return nullptr;
}
Instruction *InstCombinerImpl::visitIntToPtr(IntToPtrInst &CI) {
diff --git a/llvm/lib/Transforms/Scalar/CorrelatedValuePropagation.cpp b/llvm/lib/Transforms/Scalar/CorrelatedValuePropagation.cpp
index de3bfb57b538d3..27ca35fcdc1586 100644
--- a/llvm/lib/Transforms/Scalar/CorrelatedValuePropagation.cpp
+++ b/llvm/lib/Transforms/Scalar/CorrelatedValuePropagation.cpp
@@ -62,6 +62,7 @@ STATISTIC(NumAShrsConverted, "Number of ashr converted to lshr");
STATISTIC(NumAShrsRemoved, "Number of ashr removed");
STATISTIC(NumSRems, "Number of srem converted to urem");
STATISTIC(NumSExt, "Number of sext converted to zext");
+STATISTIC(NumSIToFP, "Number of sitofp converted to uitofp");
STATISTIC(NumSICmps, "Number of signed icmp preds simplified to unsigned");
STATISTIC(NumAnd, "Number of ands removed");
STATISTIC(NumNW, "Number of no-wrap deductions");
@@ -89,7 +90,7 @@ STATISTIC(NumSMinMax,
"Number of llvm.s{min,max} intrinsics simplified to unsigned");
STATISTIC(NumUDivURemsNarrowedExpanded,
"Number of bound udiv's/urem's expanded");
-STATISTIC(NumZExt, "Number of non-negative deductions");
+STATISTIC(NumNNeg, "Number of zext/uitofp non-negative deductions");
static Constant *getConstantAt(Value *V, Instruction *At, LazyValueInfo *LVI) {
if (Constant *C = LVI->getConstant(V, At))
@@ -1075,20 +1076,50 @@ static bool processSExt(SExtInst *SDI, LazyValueInfo *LVI) {
return true;
}
-static bool processZExt(ZExtInst *ZExt, LazyValueInfo *LVI) {
- if (ZExt->getType()->isVectorTy())
+static bool processPossibleNonNeg(Instruction *I, LazyValueInfo *LVI) {
+ assert(isa<PossiblyNonNegInst>(I) && "Requires PossiblyNonNeg instruction");
+ if (I->getType()->isVectorTy())
return false;
- if (ZExt->hasNonNeg())
+ if (I->hasNonNeg())
return false;
- const Use &Base = ZExt->getOperandUse(0);
+ const Use &Base = I->getOperandUse(0);
if (!LVI->getConstantRangeAtUse(Base, /*UndefAllowed*/ false)
.isAllNonNegative())
return false;
- ++NumZExt;
- ZExt->setNonNeg();
+ ++NumNNeg;
+ I->setNonNeg();
+
+ return true;
+}
+
+static bool processZExt(ZExtInst *ZExt, LazyValueInfo *LVI) {
+ return processPossibleNonNeg(ZExt, LVI);
+}
+
+static bool processUIToFP(UIToFPInst *UIToFP, LazyValueInfo *LVI) {
+ return processPossibleNonNeg(UIToFP, LVI);
+}
+
+static bool processSIToFP(SIToFPInst *SIToFP, LazyValueInfo *LVI) {
+ if (SIToFP->getType()->isVectorTy())
+ return false;
+
+ const Use &Base = SIToFP->getOperandUse(0);
+ if (!LVI->getConstantRangeAtUse(Base, /*UndefAllowed*/ false)
+ .isAllNonNegative())
+ return false;
+
+ ++NumSIToFP;
+ auto *UIToFP = CastInst::Create(Instruction::UIToFP, Base, SIToFP->getType(),
+ "", SIToFP->getIterator());
+ UIToFP->takeName(SIToFP);
+ UIToFP->setDebugLoc(SIToFP->getDebugLoc());
+ UIToFP->setNonNeg();
+ SIToFP->replaceAllUsesWith(UIToFP);
+ SIToFP->eraseFromParent();
return true;
}
@@ -1200,6 +1231,12 @@ static bool runImpl(Function &F, LazyValueInfo *LVI, DominatorTree *DT,
case Instruction::ZExt:
BBChanged |= processZExt(cast<ZExtInst>(&II), LVI);
break;
+ case Instruction::UIToFP:
+ BBChanged |= processUIToFP(cast<UIToFPInst>(&II), LVI);
+ break;
+ case Instruction::SIToFP:
+ BBChanged |= processSIToFP(cast<SIToFPInst>(&II), LVI);
+ break;
case Instruction::Add:
case Instruction::Sub:
case Instruction::Mul:
diff --git a/llvm/lib/Transforms/Utils/SCCPSolver.cpp b/llvm/lib/Transforms/Utils/SCCPSolver.cpp
index 38fc7763c5b204..039dfc45337552 100644
--- a/llvm/lib/Transforms/Utils/SCCPSolver.cpp
+++ b/llvm/lib/Transforms/Utils/SCCPSolver.cpp
@@ -19,6 +19,7 @@
#include "llvm/Analysis/ValueLatticeUtils.h"
#include "llvm/Analysis/ValueTracking.h"
#include "llvm/IR/InstVisitor.h"
+#include "llvm/IR/InstrTypes.h"
#include "llvm/Support/Casting.h"
#include "llvm/Support/Debug.h"
#include "llvm/Support/ErrorHandling.h"
@@ -140,7 +141,7 @@ static bool refineInstruction(SCCPSolver &Solver,
Changed = true;
}
}
- } else if (isa<ZExtInst>(Inst) && !Inst.hasNonNeg()) {
+ } else if (isa<PossiblyNonNegInst>(Inst) && !Inst.hasNonNeg()) {
auto Range = GetRange(Inst.getOperand(0));
if (Range.isAllNonNegative()) {
Inst.setNonNeg();
@@ -170,14 +171,16 @@ static bool replaceSignedInst(SCCPSolver &Solver,
Instruction *NewInst = nullptr;
switch (Inst.getOpcode()) {
- // Note: We do not fold sitofp -> uitofp here because that could be more
- // expensive in codegen and may not be reversible in the backend.
+ case Instruction::SIToFP:
case Instruction::SExt: {
- // If the source value is not negative, this is a zext.
+ // If the source value is not negative, this is a zext/uitofp.
Value *Op0 = Inst.getOperand(0);
if (InsertedValues.count(Op0) || !isNonNegative(Op0))
return false;
- NewInst = new ZExtInst(Op0, Inst.getType(), "", Inst.getIterator());
+ NewInst = CastInst::Create(Inst.getOpcode() == Instruction::SExt
+ ? Instruction::ZExt
+ : Instruction::UIToFP,
+ Op0, Inst.getType(), "", Inst.getIterator());
NewInst->setNonNeg();
break;
}
diff --git a/llvm/test/Transforms/CorrelatedValuePropagation/sitofp.ll b/llvm/test/Transforms/CorrelatedValuePropagation/sitofp.ll
new file mode 100644
index 00000000000000..db8d0a8b80e3bc
--- /dev/null
+++ b/llvm/test/Transforms/CorrelatedValuePropagation/sitofp.ll
@@ -0,0 +1,203 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py
+; RUN: opt < %s -passes=correlated-propagation -S | FileCheck %s
+
+declare void @use.f32(float)
+
+define void @test1_fptosi(i32 %n) {
+; CHECK-LABEL: @test1_fptosi(
+; CHECK-NEXT: entry:
+; CHECK-NEXT: br label [[FOR_COND:%.*]]
+; CHECK: for.cond:
+; CHECK-NEXT: [[A:%.*]] = phi i32 [ [[N:%.*]], [[ENTRY:%.*]] ], [ [[EXT:%.*]], [[FOR_BODY:%.*]] ]
+; CHECK-NEXT: [[CMP:%.*]] = icmp sgt i32 [[A]], -1
+; CHECK-NEXT: br i1 [[CMP]], label [[FOR_BODY]], label [[FOR_END:%.*]]
+; CHECK: for.body:
+; CHECK-NEXT: [[EXT_WIDE:%.*]] = uitofp nneg i32 [[A]] to float
+; CHECK-NEXT: call void @use.f32(float [[EXT_WIDE]])
+; CHECK-NEXT: [[EXT]] = fptosi float [[EXT_WIDE]] to i32
+; CHECK-NEXT: br label [[FOR_COND]]
+; CHECK: for.end:
+; CHECK-NEXT: ret void
+;
+entry:
+ br label %for.cond
+
+for.cond: ; preds = %for.body, %entry
+ %a = phi i32 [ %n, %entry ], [ %ext, %for.body ]
+ %cmp = icmp sgt i32 %a, -1
+ br i1 %cmp, label %for.body, label %for.end
+
+for.body: ; preds = %for.cond
+ %ext.wide = sitofp i32 %a to float
+ call void @use.f32(float %ext.wide)
+ %ext = fptosi float %ext.wide to i32
+ br label %for.cond
+
+for.end: ; preds = %for.cond
+ ret void
+}
+
+define void @test1_fptoui(i32 %n) {
+; CHECK-LABEL: @test1_fptoui(
+; CHECK-NEXT: entry:
+; CHECK-NEXT: br label [[FOR_COND:%.*]]
+; CHECK: for.cond:
+; CHECK-NEXT: [[A:%.*]] = phi i32 [ [[N:%.*]], [[ENTRY:%.*]] ], [ [[EXT:%.*]], [[FOR_BODY:%.*]] ]
+; CHECK-NEXT: [[CMP:%.*]] = icmp sgt i32 [[A]], -1
+; CHECK-NEXT: br i1 [[CMP]], label [[FOR_BODY]], label [[FOR_END:%.*]]
+; CHECK: for.body:
+; CHECK-NEXT: [[EXT_WIDE:%.*]] = uitofp nneg i32 [[A]] to float
+; CHECK-NEXT: call void @use.f32(float [[EXT_WIDE]])
+; CHECK-NEXT: [[EXT]] = fptoui float [[EXT_WIDE]] to i32
+; CHECK-NEXT: br label [[FOR_COND]]
+; CHECK: for.end:
+; CHECK-NEXT: ret void
+;
+entry:
+ br label %for.cond
+
+for.cond: ; preds = %for.body, %entry
+ %a = phi i32 [ %n, %entry ], [ %ext, %for.body ]
+ %cmp = icmp sgt i32 %a, -1
+ br i1 %cmp, label %for.body, label %for.end
+
+for.body: ; preds = %for.cond
+ %ext.wide = sitofp i32 %a to float
+ call void @use.f32(float %ext.wide)
+ %ext = fptoui float %ext.wide to i32
+ br label %for.cond
+
+for.end: ; preds = %for.cond
+ ret void
+}
+
+;; Negative test to show transform doesn't happen unless n >= 0.
+define void @test2(i32 %n) {
+; CHECK-LABEL: @test2(
+; CHECK-NEXT: entry:
+; CHECK-NEXT: br label [[FOR_COND:%.*]]
+; CHECK: for.cond:
+; CHECK-NEXT: [[A:%.*]] = phi i32 [ [[N:%.*]], [[ENTRY:%.*]] ], [ [[EXT:%.*]], [[FOR_BODY:%.*]] ]
+; CHECK-NEXT: [[CMP:%.*]] = icmp sgt i32 [[A]], -2
+; CHECK-NEXT: br i1 [[CMP]], label [[FOR_BODY]], label [[FOR_END:%.*]]
+; CHECK: for.body:
+; CHECK-NEXT: [[EXT_WIDE:%.*]] = sitofp i32 [[A]] to float
+; CHECK-NEXT: call void @use.f32(float [[EXT_WIDE]])
+; CHECK-NEXT: [[EXT]] = fptosi float [[EXT_WIDE]] to i32
+; CHECK-NEXT: br label [[FOR_COND]]
+; CHECK: for.end:
+; CHECK-NEXT: ret void
+;
+entry:
+ br label %for.cond
+
+for.cond: ; preds = %for.body, %entry
+ %a = phi i32 [ %n, %entry ], [ %ext, %for.body ]
+ %cmp = icmp sgt i32 %a, -2
+ br i1 %cmp, label %for.body, label %for.end
+
+for.body: ; preds = %for.cond
+ %ext.wide = sitofp i32 %a to float
+ call void @use.f32(float %ext.wide)
+ %ext = fptosi float %ext.wide to i32
+ br label %for.cond
+
+for.end: ; preds = %for.cond
+ ret void
+}
+
+;; Non looping test case.
+define void @test3(i32 %n) {
+; CHECK-LABEL: @test3(
+; CHECK-NEXT: entry:
+; CHECK-NEXT: [[CMP:%.*]] = icmp sgt i32 [[N:%.*]], -1
+; CHECK-NEXT: br i1 [[CMP]], label [[BB:%.*]], label [[EXIT:%.*]]
+; CHECK: bb:
+; CHECK-NEXT: [[EXT_WIDE:%.*]] = uitofp nneg i32 [[N]] to float
+; CHECK-NEXT: call void @use.f32(float [[EXT_WIDE]])
+; CHECK-NEXT: br label [[EXIT]]
+; CHECK: exit:
+; CHECK-NEXT: ret void
+;
+entry:
+ %cmp = icmp sgt i32 %n, -1
+ br i1 %cmp, label %bb, label %exit
+
+bb:
+ %ext.wide = sitofp i32 %n to float
+ call void @use.f32(float...
[truncated]
|
goldsteinn
changed the title
uitofp nneg
|
I looked into adding support in |
You can test this locally with the following command:git-clang-format --diff d3e77f5408fded2b4bb70f51d6d9e52684badc92 2cd2fbc5ab6e237b5cf17750f8cdb44e5c30ab41 -- llvm/lib/Transforms/Scalar/CorrelatedValuePropagation.cpp llvm/lib/Transforms/Utils/SCCPSolver.cppView the diff from clang-format here.diff --git a/llvm/lib/Transforms/Scalar/CorrelatedValuePropagation.cpp b/llvm/lib/Transforms/Scalar/CorrelatedValuePropagation.cpp
index 50b5fdb567..aaaa3a5be2 100644
--- a/llvm/lib/Transforms/Scalar/CorrelatedValuePropagation.cpp
+++ b/llvm/lib/Transforms/Scalar/CorrelatedValuePropagation.cpp
@@ -62,7 +62,7 @@ STATISTIC(NumAShrsConverted, "Number of ashr converted to lshr");
STATISTIC(NumAShrsRemoved, "Number of ashr removed");
STATISTIC(NumSRems, "Number of srem converted to urem");
STATISTIC(NumSExt, "Number of sext converted to zext");
-STATISTIC(NumSIToFP, "Number of sitofp converted to uitofp");
+STATISTIC(NumSIToFP, "Number of sitofp converted to uitofp");
STATISTIC(NumSICmps, "Number of signed icmp preds simplified to unsigned");
STATISTIC(NumAnd, "Number of ands removed");
STATISTIC(NumNW, "Number of no-wrap deductions");
|
goldsteinn
force-pushed
the
goldsteinn/uitofp-nneg-folds-and-analysis
branch
from
94808c5 to
a0c0a96
Compare
|
rebased |
|
Can you please split this into separate patches? Or at least split out the InstCombine part of it? |
See instcombine part: #88299 |
goldsteinn
changed the title
uitofp nneguitofp nneg
goldsteinn
force-pushed
the
goldsteinn/uitofp-nneg-folds-and-analysis
branch
from
a0c0a96 to
35e2a2f
Compare
goldsteinn
force-pushed
the
goldsteinn/uitofp-nneg-folds-and-analysis
branch
from
35e2a2f to
760442d
Compare
|
rebase/ping |
nikic
reviewed
May 5, 2024
|
|
||
| for.end: ; preds = %for.cond | ||
| ret void | ||
| } |
There was a problem hiding this comment.
These tests look unnecessarily complicated -- can we just assert that %a is non-negative? Or at least drop the loop around it?
| ret void | ||
| } | ||
|
|
||
| ;; Non looping test case. |
There was a problem hiding this comment.
Okay, I see you have a test without loop here -- but what's the point of the looping tests?
There was a problem hiding this comment.
The the tests are really just copy-paste of what we have for zext + nneg, didn't put much thought beyond that.
There was a problem hiding this comment.
Then please reduce the tests to the essentials... You only need the tests below here and can drop the ones above.
goldsteinn
force-pushed
the
goldsteinn/uitofp-nneg-folds-and-analysis
branch
from
760442d to
13af0e0
Compare
…tofp` -> `uitofp nneg`; NFC
Similiar to the `InstCombine` changes, just furthering the scope of the canonicalization/`uitofp nneg` support
Similiar to the `InstCombine` changes, just furthering the support of the `uitofp nneg` support.
goldsteinn
force-pushed
the
goldsteinn/uitofp-nneg-folds-and-analysis
branch
from
13af0e0 to
2cd2fbc
Compare
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
nnegflag touitofpand convertingsitofp->uitofp nneg; NFCsitofp->uitofp nnegand addnnegflag touitofpnnegflag touitofpif its operand is non-negative