maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

Separate more constant factors of parameters 

So far we separated constant factors from multiplications, however,
  only when they are at the outermost level of a parameter SCEV. Now,
  we also separate constant factors from the parameter SCEV if the
  outermost expression is a SCEVAddRecExpr. With the changes to the
  SCEVAffinator we can now improve the extractConstantFactor(...)
  function at will without worrying about any other code part. Thus,
  if needed we can implement a more comprehensive
  extractConstantFactor(...) function that will traverse the SCEV
  instead of looking only at the outermost level.

  Four test cases were affected. One did not change much and the other
  three were simplified.

llvm-svn: 260859
This commit is contained in:
Johannes Doerfert 2016-02-14 22:30:56 +00:00
parent ca741cef3e
commit 965edde695
7 changed files with 59 additions and 48 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
polly/include/polly/Support/SCEVValidator.h
View File

@ -19,6 +19,7 @@
namespace llvm {
class Region;
class SCEV;
class SCEVConstant;
class ScalarEvolution;
class Value;
class Loop;
@ -67,7 +68,7 @@ getParamsInAffineExpr(const llvm::Region *R, const llvm::SCEV *Expression,
/// @param SE The ScalarEvolution analysis to create new SCEVs.
///
/// @returns The constant factor in @p M and the rest of @p M.
std::pair<const llvm::SCEV *, const llvm::SCEV *>
std::pair<const llvm::SCEVConstant *, const llvm::SCEV *>
extractConstantFactor(const llvm::SCEV *M, llvm::ScalarEvolution &SE);
}

25
polly/lib/Support/SCEVAffinator.cpp
View File

@ -158,6 +158,10 @@ __isl_give isl_pw_aff *SCEVAffinator::visit(const SCEV *Expr) {
if (PWA)
return isl_pw_aff_copy(PWA);
auto ConstantAndLeftOverPair = extractConstantFactor(Expr, *S->getSE());
auto *Factor = ConstantAndLeftOverPair.first;
Expr = ConstantAndLeftOverPair.second;
// In case the scev is a valid parameter, we do not further analyze this
// expression, but create a new parameter in the isl_pw_aff. This allows us
// to treat subexpressions that we cannot translate into an piecewise affine
@ -171,18 +175,17 @@ __isl_give isl_pw_aff *SCEVAffinator::visit(const SCEV *Expr) {
Affine = isl_aff_add_coefficient_si(Affine, isl_dim_param, 0, 1);
PWA = isl_pw_aff_alloc(Domain, Affine);
CachedExpressions[Key] = PWA;
return isl_pw_aff_copy(PWA);
} else {
PWA = SCEVVisitor<SCEVAffinator, isl_pw_aff *>::visit(Expr);
}
PWA = SCEVVisitor<SCEVAffinator, isl_pw_aff *>::visit(Expr);
PWA = isl_pw_aff_mul(visitConstant(Factor), PWA);
// For compile time reasons we need to simplify the PWA before we cache and
// return it.
PWA = isl_pw_aff_coalesce(PWA);
CachedExpressions[Key] = PWA;
return isl_pw_aff_copy(PWA);
CachedExpressions[Key] = isl_pw_aff_copy(PWA);
return PWA;
}
__isl_give isl_pw_aff *SCEVAffinator::visitConstant(const SCEVConstant *Expr) {
@ -235,15 +238,7 @@ __isl_give isl_pw_aff *SCEVAffinator::visitAddExpr(const SCEVAddExpr *Expr) {
}
__isl_give isl_pw_aff *SCEVAffinator::visitMulExpr(const SCEVMulExpr *Expr) {
// Divide Expr into a constant part and the rest. Then visit both and multiply
// the result to obtain the representation for Expr. While the second part of
// ConstantAndLeftOverPair might still be a SCEVMulExpr we will not get to
// this point again. The reason is that if it is a multiplication it consists
// only of parameters and we will stop in the visit(const SCEV *) function and
// return the isl_pw_aff for that parameter.
auto ConstantAndLeftOverPair = extractConstantFactor(Expr, *S->getSE());
return isl_pw_aff_mul(visit(ConstantAndLeftOverPair.first),
visit(ConstantAndLeftOverPair.second));
llvm_unreachable("SCEVMulExpr should not be reached");
}
__isl_give isl_pw_aff *SCEVAffinator::visitUDivExpr(const SCEVUDivExpr *Expr) {

30
polly/lib/Support/SCEVValidator.cpp
View File

@ -640,19 +640,35 @@ std::vector<const SCEV *> getParamsInAffineExpr(const Region *R,
return Result.getParameters();
}
std::pair<const SCEV *, const SCEV *>
std::pair<const SCEVConstant *, const SCEV *>
extractConstantFactor(const SCEV *S, ScalarEvolution &SE) {
const SCEV *LeftOver = SE.getConstant(S->getType(), 1);
const SCEV *ConstPart = SE.getConstant(S->getType(), 1);
auto *LeftOver = SE.getConstant(S->getType(), 1);
auto *ConstPart = cast<SCEVConstant>(SE.getConstant(S->getType(), 1));
const SCEVMulExpr *M = dyn_cast<SCEVMulExpr>(S);
if (!M)
if (auto *Constant = dyn_cast<SCEVConstant>(S))
return std::make_pair(Constant, LeftOver);
auto *AddRec = dyn_cast<SCEVAddRecExpr>(S);
if (AddRec) {
auto *StartExpr = AddRec->getStart();
if (StartExpr->isZero()) {
auto StepPair = extractConstantFactor(AddRec->getStepRecurrence(SE), SE);
auto *LeftOverAddRec =
SE.getAddRecExpr(StartExpr, StepPair.second, AddRec->getLoop(),
AddRec->getNoWrapFlags());
return std::make_pair(StepPair.first, LeftOverAddRec);
}
return std::make_pair(ConstPart, S);
}
auto *Mul = dyn_cast<SCEVMulExpr>(S);
if (!Mul)
return std::make_pair(ConstPart, S);
for (const SCEV *Op : M->operands())
for (auto *Op : Mul->operands())
if (isa<SCEVConstant>(Op))
ConstPart = SE.getMulExpr(ConstPart, Op);
ConstPart = cast<SCEVConstant>(SE.getMulExpr(ConstPart, Op));
else
LeftOver = SE.getMulExpr(LeftOver, Op);

4
polly/test/Isl/CodeGen/OpenMP/reference-preceeding-loop.ll
View File

@ -8,9 +8,9 @@
; AST: {
; AST-NEXT: #pragma simd
; AST-NEXT: #pragma omp parallel for
; AST-NEXT: for (int c0 = 0; c0 < p_0 + symbol; c0 += 1)
; AST-NEXT: for (int c0 = 0; c0 < -p_0 + symbol; c0 += 1)
; AST-NEXT: Stmt_while_body(c0);
; AST-NEXT: if (p_0 + symbol <= 0)
; AST-NEXT: if (p_0 >= symbol)
; AST-NEXT: Stmt_while_body(0);
; AST-NEXT: }

31
polly/test/ScopInfo/NonAffine/non-affine-loop-condition-dependent-access_2.ll
View File

@ -21,15 +21,14 @@
; INNERMOST-NEXT: Invariant Accesses: {
; INNERMOST-NEXT: }
; INNERMOST-NEXT: Context:
; INNERMOST-NEXT: [p_0, p_1, p_2, p_3] -> { : 0 <= p_0 <= 1048576 and 0 <= p_1 <= 4096 and 0 <= p_2 <= 4096 and 0 <= p_3 <= 4194304 }
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { : 0 <= p_0 <= 1048576 and 0 <= p_1 <= 1024 and 0 <= p_2 <= 1024 }
; INNERMOST-NEXT: Assumed Context:
; INNERMOST-NEXT: [p_0, p_1, p_2, p_3] -> { : }
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { : }
; INNERMOST-NEXT: Boundary Context:
; INNERMOST-NEXT: [p_0, p_1, p_2, p_3] -> { : }
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { : }
; INNERMOST-NEXT: p0: {0,+,{0,+,1}<nuw><nsw><%bb11>}<nuw><nsw><%bb13>
; INNERMOST-NEXT: p1: {0,+,4}<nuw><nsw><%bb11>
; INNERMOST-NEXT: p2: {0,+,4}<nuw><nsw><%bb13>
; INNERMOST-NEXT: p3: {0,+,{0,+,4}<nuw><nsw><%bb11>}<%bb13>
; INNERMOST-NEXT: p1: {0,+,1}<nuw><nsw><%bb11>
; INNERMOST-NEXT: p2: {0,+,1}<nuw><nsw><%bb13>
; INNERMOST-NEXT: Arrays {
; INNERMOST-NEXT: i32 MemRef_A[*]; // Element size 4
; INNERMOST-NEXT: i64 MemRef_indvars_iv_next6; // Element size 8
@ -45,26 +44,26 @@
; INNERMOST-NEXT: Statements {
; INNERMOST-NEXT: Stmt_bb16
; INNERMOST-NEXT: Domain :=
; INNERMOST-NEXT: [p_0, p_1, p_2, p_3] -> { Stmt_bb16[i0] : 0 <= i0 <= 1023 - p_0 };
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb16[i0] : 0 <= i0 <= 1023 - p_0 };
; INNERMOST-NEXT: Schedule :=
; INNERMOST-NEXT: [p_0, p_1, p_2, p_3] -> { Stmt_bb16[i0] -> [0, i0] };
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb16[i0] -> [0, i0] };
; INNERMOST-NEXT: ReadAccess := [Reduction Type: NONE] [Scalar: 0]
; INNERMOST-NEXT: [p_0, p_1, p_2, p_3] -> { Stmt_bb16[i0] -> MemRef_A[o0] : 4o0 = p_1 };
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb16[i0] -> MemRef_A[p_1] };
; INNERMOST-NEXT: ReadAccess := [Reduction Type: NONE] [Scalar: 0]
; INNERMOST-NEXT: [p_0, p_1, p_2, p_3] -> { Stmt_bb16[i0] -> MemRef_A[o0] : 4o0 = p_2 };
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb16[i0] -> MemRef_A[p_2] };
; INNERMOST-NEXT: ReadAccess := [Reduction Type: +] [Scalar: 0]
; INNERMOST-NEXT: [p_0, p_1, p_2, p_3] -> { Stmt_bb16[i0] -> MemRef_A[o0] : 4o0 = p_3 + 4i0 };
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb16[i0] -> MemRef_A[p_0 + i0] };
; INNERMOST-NEXT: MustWriteAccess := [Reduction Type: +] [Scalar: 0]
; INNERMOST-NEXT: [p_0, p_1, p_2, p_3] -> { Stmt_bb16[i0] -> MemRef_A[o0] : 4o0 = p_3 + 4i0 };
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb16[i0] -> MemRef_A[p_0 + i0] };
; INNERMOST-NEXT: Stmt_bb26
; INNERMOST-NEXT: Domain :=
; INNERMOST-NEXT: [p_0, p_1, p_2, p_3] -> { Stmt_bb26[] : p_0 <= 1024 };
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb26[] : p_0 <= 1024 };
; INNERMOST-NEXT: Schedule :=
; INNERMOST-NEXT: [p_0, p_1, p_2, p_3] -> { Stmt_bb26[] -> [1, 0] };
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb26[] -> [1, 0] };
; INNERMOST-NEXT: MustWriteAccess := [Reduction Type: NONE] [Scalar: 1]
; INNERMOST-NEXT: [p_0, p_1, p_2, p_3] -> { Stmt_bb26[] -> MemRef_indvars_iv_next6[] };
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb26[] -> MemRef_indvars_iv_next6[] };
; INNERMOST-NEXT: MustWriteAccess := [Reduction Type: NONE] [Scalar: 1]
; INNERMOST-NEXT: [p_0, p_1, p_2, p_3] -> { Stmt_bb26[] -> MemRef_indvars_iv_next4[] };
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb26[] -> MemRef_indvars_iv_next4[] };
; INNERMOST-NEXT: }
; ALL: Function: f

10
polly/test/ScopInfo/NonAffine/non-affine-loop-condition-dependent-access_3.ll
View File

@ -20,14 +20,14 @@
; INNERMOST-NEXT: Invariant Accesses: {
; INNERMOST-NEXT: }
; INNERMOST-NEXT: Context:
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { : 0 <= p_0 <= 2147483647 and 0 <= p_1 <= 4096 and 0 <= p_2 <= 4096 }
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { : 0 <= p_0 <= 2147483647 and 0 <= p_1 <= 1024 and 0 <= p_2 <= 1024 }
; INNERMOST-NEXT: Assumed Context:
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { : }
; INNERMOST-NEXT: Boundary Context:
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { : }
; INNERMOST-NEXT: p0: {0,+,{0,+,1}<nuw><nsw><%bb11>}<nuw><nsw><%bb13>
; INNERMOST-NEXT: p1: {0,+,4}<nuw><nsw><%bb11>
; INNERMOST-NEXT: p2: {0,+,4}<nuw><nsw><%bb13>
; INNERMOST-NEXT: p1: {0,+,1}<nuw><nsw><%bb11>
; INNERMOST-NEXT: p2: {0,+,1}<nuw><nsw><%bb13>
; INNERMOST-NEXT: Arrays {
; INNERMOST-NEXT: i32 MemRef_A[*]; // Element size 4
; INNERMOST-NEXT: i64 MemRef_indvars_iv_next6; // Element size 8
@ -47,9 +47,9 @@
; INNERMOST-NEXT: Schedule :=
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb16[i0] -> [0, i0] };
; INNERMOST-NEXT: ReadAccess := [Reduction Type: NONE] [Scalar: 0]
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb16[i0] -> MemRef_A[o0] : 4o0 = p_1 };
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb16[i0] -> MemRef_A[p_1] };
; INNERMOST-NEXT: ReadAccess := [Reduction Type: NONE] [Scalar: 0]
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb16[i0] -> MemRef_A[o0] : 4o0 = p_2 };
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb16[i0] -> MemRef_A[p_2] };
; INNERMOST-NEXT: ReadAccess := [Reduction Type: +] [Scalar: 0]
; INNERMOST-NEXT: [p_0, p_1, p_2] -> { Stmt_bb16[i0] -> MemRef_A[i0] };
; INNERMOST-NEXT: MustWriteAccess := [Reduction Type: +] [Scalar: 0]

4
polly/test/ScopInfo/constant_start_integer.ll
View File

@ -15,9 +15,9 @@ target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f3
; CHECK-NOT: p1
; CHECK: ReadAccess
; CHECK: [p_0] -> { Stmt_for_body3[i0] -> MemRef_input[o0] : 4o0 = 4 + p_0 + 4i0 };
; CHECK: [p_0] -> { Stmt_for_body3[i0] -> MemRef_input[1 + 64p_0 + i0] };
; CHECK: MustWriteAccess
; CHECK: [p_0] -> { Stmt_for_body3[i0] -> MemRef_input[o0] : 4o0 = p_0 + 4i0 };
; CHECK: [p_0] -> { Stmt_for_body3[i0] -> MemRef_input[64p_0 + i0] };
define void @foo(float* nocapture %input) {
entry: