From 5b9ff8b667f932cb742c4b60f16a7ac39d8707dc Mon Sep 17 00:00:00 2001 From: Johannes Doerfert Date: Thu, 10 Sep 2015 13:00:06 +0000 Subject: [PATCH] Replace ScalarEvolution based domain generation This patch replaces the last legacy part of the domain generation, namely the ScalarEvolution part that was used to obtain loop bounds. We now iterate over the loops in the region and propagate the back edge condition to the header blocks. Afterwards we propagate the new information once through the whole region. In this process we simply ignore unbounded parts of the domain and thereby assume the absence of infinite loops. + This patch already identified a couple of broken unit tests we had for years. + We allow more loops already and the step to multiple exit and multiple back edges is minimal. + It allows to model the overflow checks properly as we actually visit every block in the SCoP and know where which condition is evaluated. - It is currently not compatible with modulo constraints in the domain. Differential Revision: http://reviews.llvm.org/D12499 llvm-svn: 247279 --- polly/include/polly/ScopInfo.h | 20 +- polly/lib/Analysis/ScopInfo.cpp | 401 +++++++++++++----- .../loop-body-references-outer-values-3.ll | 6 +- .../OpenMP/reference-preceeding-loop.ll | 22 +- ...o-parallel-loops-reference-outer-indvar.ll | 11 +- polly/test/Isl/CodeGen/phi_scalar_simple_1.ll | 42 +- polly/test/Isl/CodeGen/two-scops-in-row.ll | 5 +- polly/test/ScopInfo/NonAffine/div_backedge.ll | 42 ++ polly/test/ScopInfo/NonAffine/div_domain.ll | 54 +++ .../ScopInfo/NonAffine/modulo_backedge.ll | 44 ++ .../test/ScopInfo/NonAffine/modulo_domain.ll | 56 +++ ...ffine-loop-condition-dependent-access_2.ll | 27 +- ...ffine-loop-condition-dependent-access_3.ll | 2 +- polly/test/ScopInfo/isl_trip_count_01.ll | 2 +- polly/test/ScopInfo/loop_affine_bound_0.ll | 2 - polly/test/ScopInfo/loop_affine_bound_1.ll | 10 +- polly/test/ScopInfo/loop_affine_bound_2.ll | 2 - polly/test/ScopInfo/non_affine_region_1.ll | 6 +- .../test/ScopInfo/pointer-type-expressions.ll | 6 +- .../redundant_parameter_constraint.ll | 43 ++ .../two-loops-right-after-each-other.ll | 16 +- polly/test/ScopInfo/unsigned-condition.ll | 6 +- 22 files changed, 630 insertions(+), 195 deletions(-) create mode 100644 polly/test/ScopInfo/NonAffine/div_backedge.ll create mode 100644 polly/test/ScopInfo/NonAffine/div_domain.ll create mode 100644 polly/test/ScopInfo/NonAffine/modulo_backedge.ll create mode 100644 polly/test/ScopInfo/NonAffine/modulo_domain.ll create mode 100644 polly/test/ScopInfo/redundant_parameter_constraint.ll diff --git a/polly/include/polly/ScopInfo.h b/polly/include/polly/ScopInfo.h index a122d98c9db1..fd5c81264446 100644 --- a/polly/include/polly/ScopInfo.h +++ b/polly/include/polly/ScopInfo.h @@ -659,9 +659,6 @@ private: /// Build the statement. //@{ - void addConditionsToDomain(TempScop &tempScop, const Region &CurRegion); - void addLoopBoundsToDomain(TempScop &tempScop); - void addLoopTripCountToDomain(const Loop *L); void buildDomain(TempScop &tempScop, const Region &CurRegion); /// @brief Create the accesses for instructions in @p Block. @@ -1018,6 +1015,14 @@ private: AliasAnalysis &AA, DominatorTree &DT, isl_ctx *ctx); + /// @brief Add loop carried constraints to the header blocks of loops. + /// + /// @param LI The LoopInfo analysis. + /// @param SD The ScopDetection analysis to identify non-affine sub-regions. + /// @param DT The dominator tree of the current function. + void addLoopBoundsToHeaderDomains(LoopInfo &LI, ScopDetection &SD, + DominatorTree &DT); + /// @brief Compute the branching constraints for each basic block in @p R. /// /// @param R The region we currently build branching conditions for. @@ -1027,6 +1032,15 @@ private: void buildDomainsWithBranchConstraints(Region *R, LoopInfo &LI, ScopDetection &SD, DominatorTree &DT); + /// @brief Propagate the domain constraints through the region @p R. + /// + /// @param R The region we currently build branching conditions for. + /// @param LI The LoopInfo analysis to obtain the number of iterators. + /// @param SD The ScopDetection analysis to identify non-affine sub-regions. + /// @param DT The dominator tree of the current function. + void propagateDomainConstraints(Region *R, LoopInfo &LI, ScopDetection &SD, + DominatorTree &DT); + /// @brief Compute the domain for each basic block in @p R. /// /// @param R The region we currently traverse. diff --git a/polly/lib/Analysis/ScopInfo.cpp b/polly/lib/Analysis/ScopInfo.cpp index f3104a992ed4..e27026fb18eb 100644 --- a/polly/lib/Analysis/ScopInfo.cpp +++ b/polly/lib/Analysis/ScopInfo.cpp @@ -769,6 +769,68 @@ void ScopStmt::realignParams() { Domain = isl_set_align_params(Domain, Parent.getParamSpace()); } +/// @brief Add @p BSet to the set @p User if @p BSet is bounded. +static isl_stat collectBoundedParts(__isl_take isl_basic_set *BSet, + void *User) { + isl_set **BoundedParts = static_cast(User); + if (isl_basic_set_is_bounded(BSet)) + *BoundedParts = isl_set_union(*BoundedParts, isl_set_from_basic_set(BSet)); + else + isl_basic_set_free(BSet); + return isl_stat_ok; +} + +/// @brief Return the bounded parts of @p S. +static __isl_give isl_set *collectBoundedParts(__isl_take isl_set *S) { + isl_set *BoundedParts = isl_set_empty(isl_set_get_space(S)); + isl_set_foreach_basic_set(S, collectBoundedParts, &BoundedParts); + isl_set_free(S); + return BoundedParts; +} + +/// @brief Compute the (un)bounded parts of @p S wrt. to dimension @p Dim. +/// +/// @returns A separation of @p S into first an unbounded then a bounded subset, +/// both with regards to the dimension @p Dim. +static std::pair<__isl_give isl_set *, __isl_give isl_set *> +partitionSetParts(__isl_take isl_set *S, unsigned Dim) { + + for (unsigned u = 0, e = isl_set_n_dim(S); u < e; u++) + S = isl_set_lower_bound_si(S, isl_dim_set, u, 0); + + unsigned NumDimsS = isl_set_n_dim(S); + isl_set *OnlyDimS = isl_set_copy(S); + + // Remove dimensions that are greater than Dim as they are not interesting. + assert(NumDimsS >= Dim + 1); + OnlyDimS = + isl_set_project_out(OnlyDimS, isl_dim_set, Dim + 1, NumDimsS - Dim - 1); + + // Create artificial parametric upper bounds for dimensions smaller than Dim + // as we are not interested in them. + OnlyDimS = isl_set_insert_dims(OnlyDimS, isl_dim_param, 0, Dim); + for (unsigned u = 0; u < Dim; u++) { + isl_constraint *C = isl_inequality_alloc( + isl_local_space_from_space(isl_set_get_space(OnlyDimS))); + C = isl_constraint_set_coefficient_si(C, isl_dim_param, u, 1); + C = isl_constraint_set_coefficient_si(C, isl_dim_set, u, -1); + OnlyDimS = isl_set_add_constraint(OnlyDimS, C); + } + + // Collect all bounded parts of OnlyDimS. + isl_set *BoundedParts = collectBoundedParts(OnlyDimS); + + // Create the dimensions greater than Dim again. + BoundedParts = isl_set_insert_dims(BoundedParts, isl_dim_set, Dim + 1, + NumDimsS - Dim - 1); + + // Remove the artificial upper bound parameters again. + BoundedParts = isl_set_remove_dims(BoundedParts, isl_dim_param, 0, Dim); + + isl_set *UnboundedParts = isl_set_subtract(S, isl_set_copy(BoundedParts)); + return std::make_pair(UnboundedParts, BoundedParts); +} + static __isl_give isl_set *buildConditionSet(ICmpInst::Predicate Pred, isl_pw_aff *L, isl_pw_aff *R) { switch (Pred) { @@ -841,121 +903,12 @@ buildConditionSets(Scop &S, BranchInst *BI, Loop *L, __isl_keep isl_set *Domain, isl_set_intersect(AlternativeCondSet, isl_set_copy(Domain)))); } -void ScopStmt::addLoopTripCountToDomain(const Loop *L) { - - int RelativeLoopDimension = getParent()->getRelativeLoopDepth(L); - assert(RelativeLoopDimension >= 0 && - "Expected relative loop depth of L to be non-negative"); - unsigned loopDimension = RelativeLoopDimension; - - ScalarEvolution *SE = getParent()->getSE(); - isl_space *DomSpace = isl_set_get_space(Domain); - - isl_space *MapSpace = isl_space_map_from_set(isl_space_copy(DomSpace)); - isl_multi_aff *LoopMAff = isl_multi_aff_identity(MapSpace); - isl_aff *LoopAff = isl_multi_aff_get_aff(LoopMAff, loopDimension); - LoopAff = isl_aff_add_constant_si(LoopAff, 1); - LoopMAff = isl_multi_aff_set_aff(LoopMAff, loopDimension, LoopAff); - isl_map *TranslationMap = isl_map_from_multi_aff(LoopMAff); - - BasicBlock *ExitingBB = L->getExitingBlock(); - assert(ExitingBB && "Loop has more than one exiting block"); - - BranchInst *Term = dyn_cast(ExitingBB->getTerminator()); - assert(Term && Term->isConditional() && "Terminator is not conditional"); - - const SCEV *LHS = nullptr; - const SCEV *RHS = nullptr; - Value *Cond = Term->getCondition(); - CmpInst::Predicate Pred = CmpInst::Predicate::BAD_ICMP_PREDICATE; - - ICmpInst *CondICmpInst = dyn_cast(Cond); - ConstantInt *CondConstant = dyn_cast(Cond); - if (CondICmpInst) { - LHS = SE->getSCEVAtScope(CondICmpInst->getOperand(0), L); - RHS = SE->getSCEVAtScope(CondICmpInst->getOperand(1), L); - Pred = CondICmpInst->getPredicate(); - } else if (CondConstant) { - LHS = SE->getConstant(CondConstant); - RHS = SE->getConstant(ConstantInt::getTrue(SE->getContext())); - Pred = CmpInst::Predicate::ICMP_EQ; - } else { - llvm_unreachable("Condition is neither a ConstantInt nor a ICmpInst"); - } - - if (!L->contains(Term->getSuccessor(0))) - Pred = ICmpInst::getInversePredicate(Pred); - Comparison Comp(LHS, RHS, Pred); - - isl_pw_aff *LPWA = getPwAff(Comp.getLHS()); - isl_pw_aff *RPWA = getPwAff(Comp.getRHS()); - - isl_set *CondSet = buildConditionSet(Comp.getPred(), LPWA, RPWA); - isl_map *ForwardMap = isl_map_lex_le(isl_space_copy(DomSpace)); - for (unsigned i = 0; i < isl_set_n_dim(Domain); i++) - if (i != loopDimension) - ForwardMap = isl_map_equate(ForwardMap, isl_dim_in, i, isl_dim_out, i); - - ForwardMap = isl_map_apply_range(ForwardMap, isl_map_copy(TranslationMap)); - isl_set *CondDom = isl_set_subtract(isl_set_copy(Domain), CondSet); - isl_set *ForwardCond = isl_set_apply(CondDom, isl_map_copy(ForwardMap)); - isl_set *ForwardDomain = isl_set_apply(isl_set_copy(Domain), ForwardMap); - ForwardCond = isl_set_gist(ForwardCond, ForwardDomain); - Domain = isl_set_subtract(Domain, ForwardCond); - - isl_map_free(TranslationMap); - isl_space_free(DomSpace); -} - -void ScopStmt::addLoopBoundsToDomain(TempScop &tempScop) { - isl_space *Space; - isl_local_space *LocalSpace; - - Space = isl_set_get_space(Domain); - LocalSpace = isl_local_space_from_space(Space); - - ScalarEvolution *SE = getParent()->getSE(); - for (int i = 0, e = getNumIterators(); i != e; ++i) { - isl_aff *Zero = isl_aff_zero_on_domain(isl_local_space_copy(LocalSpace)); - isl_pw_aff *IV = - isl_pw_aff_from_aff(isl_aff_set_coefficient_si(Zero, isl_dim_in, i, 1)); - - // 0 <= IV. - isl_set *LowerBound = isl_pw_aff_nonneg_set(isl_pw_aff_copy(IV)); - Domain = isl_set_intersect(Domain, LowerBound); - - // IV <= LatchExecutions. - const Loop *L = getLoopForDimension(i); - const SCEV *LatchExecutions = SE->getBackedgeTakenCount(L); - if (!isa(LatchExecutions)) { - isl_pw_aff *UpperBound = getPwAff(LatchExecutions); - isl_set *UpperBoundSet = isl_pw_aff_le_set(IV, UpperBound); - Domain = isl_set_intersect(Domain, UpperBoundSet); - } else { - // If SCEV cannot provide a loop trip count, we compute it with ISL. If - // the domain remains unbounded, make the assumed context infeasible - // as code generation currently does not expect unbounded loops. - addLoopTripCountToDomain(L); - isl_pw_aff_free(IV); - if (!isl_set_dim_has_upper_bound(Domain, isl_dim_set, i)) - Parent.addAssumption(isl_set_empty(Parent.getParamSpace())); - } - } - - isl_local_space_free(LocalSpace); -} - void ScopStmt::buildDomain(TempScop &tempScop, const Region &CurRegion) { - isl_space *Space; isl_id *Id; - Space = isl_space_set_alloc(getIslCtx(), 0, getNumIterators()); - Id = isl_id_alloc(getIslCtx(), getBaseName(), this); - Domain = isl_set_universe(Space); - addLoopBoundsToDomain(tempScop); - Domain = isl_set_intersect(Domain, getParent()->getDomainConditions(this)); + Domain = getParent()->getDomainConditions(this); Domain = isl_set_coalesce(Domain); Domain = isl_set_set_tuple_id(Domain, Id); } @@ -1533,6 +1486,8 @@ void Scop::buildDomains(Region *R, LoopInfo &LI, ScopDetection &SD, DomainMap[EntryBB] = S; buildDomainsWithBranchConstraints(R, LI, SD, DT); + addLoopBoundsToHeaderDomains(LI, SD, DT); + propagateDomainConstraints(R, LI, SD, DT); } void Scop::buildDomainsWithBranchConstraints(Region *R, LoopInfo &LI, @@ -1639,6 +1594,232 @@ void Scop::buildDomainsWithBranchConstraints(Region *R, LoopInfo &LI, } } +/// @brief Return the domain for @p BB wrt @p DomainMap. +/// +/// This helper function will lookup @p BB in @p DomainMap but also handle the +/// case where @p BB is contained in a non-affine subregion using the region +/// tree obtained by @p RI. +static __isl_give isl_set * +getDomainForBlock(BasicBlock *BB, DenseMap &DomainMap, + RegionInfo &RI) { + auto DIt = DomainMap.find(BB); + if (DIt != DomainMap.end()) + return isl_set_copy(DIt->getSecond()); + + Region *R = RI.getRegionFor(BB); + while (R->getEntry() == BB) + R = R->getParent(); + return getDomainForBlock(R->getEntry(), DomainMap, RI); +} + +void Scop::propagateDomainConstraints(Region *R, LoopInfo &LI, + ScopDetection &SD, DominatorTree &DT) { + // Iterate over the region R and propagate the domain constrains from the + // predecessors to the current node. In contrast to the + // buildDomainsWithBranchConstraints function, this one will pull the domain + // information from the predecessors instead of pushing it to the successors. + // Additionally, we assume the domains to be already present in the domain + // map here. However, we iterate again in reverse post order so we know all + // predecessors have been visited before a block or non-affine subregion is + // visited. + + // The set of boxed loops (loops in non-affine subregions) for this SCoP. + auto &BoxedLoops = *SD.getBoxedLoops(&getRegion()); + + ReversePostOrderTraversal RTraversal(R); + for (auto *RN : RTraversal) { + + // Recurse for affine subregions but go on for basic blocks and non-affine + // subregions. + if (RN->isSubRegion()) { + Region *SubRegion = RN->getNodeAs(); + if (!SD.isNonAffineSubRegion(SubRegion, &getRegion())) { + propagateDomainConstraints(SubRegion, LI, SD, DT); + continue; + } + } + + BasicBlock *BB = getRegionNodeBasicBlock(RN); + Loop *BBLoop = getRegionNodeLoop(RN, LI); + int BBLoopDepth = getRelativeLoopDepth(BBLoop); + + isl_set *&Domain = DomainMap[BB]; + assert(Domain && "Due to reverse post order traversal of the region all " + "predecessor of the current region node should have been " + "visited and a domain for this region node should have " + "been set."); + + isl_set *PredDom = isl_set_empty(isl_set_get_space(Domain)); + for (auto *PredBB : predecessors(BB)) { + + // Skip backedges + if (DT.dominates(BB, PredBB)) + continue; + + isl_set *PredBBDom = nullptr; + + // Handle the SCoP entry block with its outside predecessors. + if (!getRegion().contains(PredBB)) + PredBBDom = isl_set_universe(isl_set_get_space(PredDom)); + + if (!PredBBDom) { + // Determine the loop depth of the predecessor and adjust its domain to + // the domain of the current block. This can mean we have to: + // o) Drop a dimension if this block is the exit of a loop, not the + // header of a new loop and the predecessor was part of the loop. + // o) Add an unconstrainted new dimension if this block is the header + // of a loop and the predecessor is not part of it. + // o) Drop the information about the innermost loop dimension when the + // predecessor and the current block are surrounded by different + // loops in the same depth. + PredBBDom = getDomainForBlock(PredBB, DomainMap, *R->getRegionInfo()); + Loop *PredBBLoop = LI.getLoopFor(PredBB); + while (BoxedLoops.count(PredBBLoop)) + PredBBLoop = PredBBLoop->getParentLoop(); + + int PredBBLoopDepth = getRelativeLoopDepth(PredBBLoop); + assert(std::abs(BBLoopDepth - PredBBLoopDepth) <= 1); + if (BBLoopDepth < PredBBLoopDepth) + PredBBDom = + isl_set_project_out(PredBBDom, isl_dim_set, PredBBLoopDepth, 1); + else if (PredBBLoopDepth < BBLoopDepth) + PredBBDom = isl_set_add_dims(PredBBDom, isl_dim_set, 1); + else if (BBLoop != PredBBLoop && BBLoopDepth >= 0) + PredBBDom = isl_set_drop_constraints_involving_dims( + PredBBDom, isl_dim_set, BBLoopDepth, 1); + } + + PredDom = isl_set_union(PredDom, PredBBDom); + } + + // Under the union of all predecessor conditions we can reach this block. + Domain = isl_set_intersect(Domain, PredDom); + } +} + +/// @brief Create a map from SetSpace -> SetSpace where the dimensions @p Dim +/// is incremented by one and all other dimensions are equal, e.g., +/// [i0, i1, i2, i3] -> [i0, i1, i2 + 1, i3] +/// if @p Dim is 2 and @p SetSpace has 4 dimensions. +static __isl_give isl_map * +createNextIterationMap(__isl_take isl_space *SetSpace, unsigned Dim) { + auto *MapSpace = isl_space_map_from_set(SetSpace); + auto *NextIterationMap = isl_map_universe(isl_space_copy(MapSpace)); + for (unsigned u = 0; u < isl_map_n_in(NextIterationMap); u++) + if (u != Dim) + NextIterationMap = + isl_map_equate(NextIterationMap, isl_dim_in, u, isl_dim_out, u); + auto *C = isl_constraint_alloc_equality(isl_local_space_from_space(MapSpace)); + C = isl_constraint_set_constant_si(C, 1); + C = isl_constraint_set_coefficient_si(C, isl_dim_in, Dim, 1); + C = isl_constraint_set_coefficient_si(C, isl_dim_out, Dim, -1); + NextIterationMap = isl_map_add_constraint(NextIterationMap, C); + return NextIterationMap; +} + +/// @brief Add @p L & all children to @p Loops if they are not in @p BoxedLoops. +static inline void +addLoopAndSubloops(Loop *L, SmallVectorImpl &Loops, + const ScopDetection::BoxedLoopsSetTy &BoxedLoops) { + if (BoxedLoops.count(L)) + return; + + Loops.push_back(L); + for (Loop *Subloop : *L) + addLoopAndSubloops(Subloop, Loops, BoxedLoops); +} + +/// @brief Add loops in @p R to @p RegionLoops if they are not in @p BoxedLoops. +static inline void +collectLoopsInRegion(Region &R, LoopInfo &LI, + SmallVector &RegionLoops, + const ScopDetection::BoxedLoopsSetTy &BoxedLoops) { + + SmallVector Loops(LI.begin(), LI.end()); + while (!Loops.empty()) { + Loop *L = Loops.pop_back_val(); + + if (R.contains(L)) + addLoopAndSubloops(L, RegionLoops, BoxedLoops); + else if (L->contains(R.getEntry())) + Loops.append(L->begin(), L->end()); + } +} + +/// @brief Create a set from @p Space with @p Dim fixed to 0. +static __isl_give isl_set * +createFirstIterationDomain(__isl_take isl_space *Space, unsigned Dim) { + auto *Domain = isl_set_universe(Space); + Domain = isl_set_fix_si(Domain, isl_dim_set, Dim, 0); + return Domain; +} + +void Scop::addLoopBoundsToHeaderDomains(LoopInfo &LI, ScopDetection &SD, + DominatorTree &DT) { + // We iterate over all loops in the SCoP, create the condition set under which + // we will take the back edge, and then apply these restrictions to the + // header. + + Region &R = getRegion(); + SmallVector RegionLoops; + collectLoopsInRegion(R, LI, RegionLoops, *SD.getBoxedLoops(&R)); + + while (!RegionLoops.empty()) { + Loop *L = RegionLoops.pop_back_val(); + int LoopDepth = getRelativeLoopDepth(L); + assert(LoopDepth >= 0 && "Loop in region should have at least depth one"); + + BasicBlock *LatchBB = L->getLoopLatch(); + assert(LatchBB && "TODO implement multiple exit loop handling"); + + isl_set *LatchBBDom = DomainMap[LatchBB]; + isl_set *BackedgeCondition = nullptr; + + BasicBlock *HeaderBB = L->getHeader(); + + BranchInst *BI = cast(LatchBB->getTerminator()); + if (BI->isUnconditional()) + BackedgeCondition = isl_set_copy(LatchBBDom); + else { + SmallVector ConditionSets; + int idx = BI->getSuccessor(0) != HeaderBB; + buildConditionSets(*this, BI, L, LatchBBDom, ConditionSets); + + // Free the non back edge condition set as we do not need it. + isl_set_free(ConditionSets[1 - idx]); + + BackedgeCondition = ConditionSets[idx]; + } + + isl_set *&HeaderBBDom = DomainMap[HeaderBB]; + isl_set *FirstIteration = + createFirstIterationDomain(isl_set_get_space(HeaderBBDom), LoopDepth); + + isl_map *NextIterationMap = + createNextIterationMap(isl_set_get_space(HeaderBBDom), LoopDepth); + + int LatchLoopDepth = getRelativeLoopDepth(LI.getLoopFor(LatchBB)); + assert(LatchLoopDepth >= LoopDepth); + BackedgeCondition = + isl_set_project_out(BackedgeCondition, isl_dim_set, LoopDepth + 1, + LatchLoopDepth - LoopDepth); + + auto Parts = partitionSetParts(BackedgeCondition, LoopDepth); + + // If a loop has an unbounded back edge condition part (here Parts.first) + // we do not want to assume the header will even be executed for the first + // iteration of an execution that will lead to an infinite loop. While it + // would not be wrong to do so, it does not seem helpful. + FirstIteration = isl_set_subtract(FirstIteration, Parts.first); + + BackedgeCondition = isl_set_apply(Parts.second, NextIterationMap); + BackedgeCondition = isl_set_union(BackedgeCondition, FirstIteration); + BackedgeCondition = isl_set_coalesce(BackedgeCondition); + + HeaderBBDom = isl_set_intersect(HeaderBBDom, BackedgeCondition); + } +} + void Scop::buildAliasChecks(AliasAnalysis &AA) { if (!PollyUseRuntimeAliasChecks) return; diff --git a/polly/test/Isl/CodeGen/OpenMP/loop-body-references-outer-values-3.ll b/polly/test/Isl/CodeGen/OpenMP/loop-body-references-outer-values-3.ll index b79503aa3e8a..3c36784be112 100644 --- a/polly/test/Isl/CodeGen/OpenMP/loop-body-references-outer-values-3.ll +++ b/polly/test/Isl/CodeGen/OpenMP/loop-body-references-outer-values-3.ll @@ -8,8 +8,9 @@ ; but %call is a parameter of the SCoP and we need to make sure its value is ; properly forwarded to the subfunction. +; AST: Stmt_for_body(0); ; AST: #pragma omp parallel for -; AST: for (int c0 = 0; c0 < cols; c0 += 1) +; AST: for (int c0 = 1; c0 < cols; c0 += 1) ; AST: Stmt_for_body(c0); ; IR: @foo_polly_subfn @@ -38,8 +39,9 @@ end: ; Another variation of this test case, now with even more of the index ; expression defined outside of the scop. +; AST: Stmt_for_body(0); ; AST: #pragma omp parallel for -; AST: for (int c0 = 0; c0 < cols; c0 += 1) +; AST: for (int c0 = 1; c0 < cols; c0 += 1) ; AST: Stmt_for_body(c0); ; IR: @bar_polly_subfn diff --git a/polly/test/Isl/CodeGen/OpenMP/reference-preceeding-loop.ll b/polly/test/Isl/CodeGen/OpenMP/reference-preceeding-loop.ll index 68086d7d80c1..c7cdf39ff635 100644 --- a/polly/test/Isl/CodeGen/OpenMP/reference-preceeding-loop.ll +++ b/polly/test/Isl/CodeGen/OpenMP/reference-preceeding-loop.ll @@ -4,21 +4,17 @@ ; - Test the case where scalar evolution references a loop that is outside ; of the scop, but does not contain the scop. -; - Test the case where two parallel subfunctions are created. - -; AST: if (symbol >= p_2 + 1) { -; AST-NEXT: #pragma simd -; AST-NEXT: #pragma omp parallel for -; AST-NEXT: for (int c0 = 0; c0 < p_0 + symbol; c0 += 1) -; AST-NEXT: Stmt_while_body(c0); -; AST-NEXT: } else -; AST-NEXT: #pragma simd -; AST-NEXT: #pragma omp parallel for -; AST-NEXT: for (int c0 = 0; c0 <= p_0 + p_2; c0 += 1) -; AST-NEXT: Stmt_while_body(c0); + +; AST: { +; AST-NEXT: Stmt_while_body(0); +; AST-NEXT: #pragma simd +; AST-NEXT: #pragma omp parallel for +; AST-NEXT: for (int c0 = 1; c0 < p_0 + symbol; c0 += 1) +; AST-NEXT: Stmt_while_body(c0); +; AST-NEXT: } ; IR: @update_model_polly_subfn -; IR: @update_model_polly_subfn_1 +; IR-NOT: @update_model_polly_subfn_1 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" diff --git a/polly/test/Isl/CodeGen/OpenMP/two-parallel-loops-reference-outer-indvar.ll b/polly/test/Isl/CodeGen/OpenMP/two-parallel-loops-reference-outer-indvar.ll index 8755af284558..d5e4106c9207 100644 --- a/polly/test/Isl/CodeGen/OpenMP/two-parallel-loops-reference-outer-indvar.ll +++ b/polly/test/Isl/CodeGen/OpenMP/two-parallel-loops-reference-outer-indvar.ll @@ -4,19 +4,14 @@ ; This test case verifies that we create correct code even if two OpenMP loops ; share common outer variables. -; AST: if (nj >= p_0 + 3) { +; AST: Stmt_for_body35(0); ; AST: #pragma simd ; AST: #pragma omp parallel for -; AST: for (int c0 = 0; c0 < nj + p_2 - 1; c0 += 1) -; AST: Stmt_for_body35(c0); -; AST: } else -; AST: #pragma simd -; AST: #pragma omp parallel for -; AST: for (int c0 = 0; c0 <= p_0 + p_2; c0 += 1) +; AST: for (int c0 = 1; c0 < -p_0 + nj - 1; c0 += 1) ; AST: Stmt_for_body35(c0); ; IR: @foo_polly_subfn -; IR: @foo_polly_subfn_1 +; IR-NOT: @foo_polly_subfn_1 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" diff --git a/polly/test/Isl/CodeGen/phi_scalar_simple_1.ll b/polly/test/Isl/CodeGen/phi_scalar_simple_1.ll index 13cab99dff02..25290d9c2180 100644 --- a/polly/test/Isl/CodeGen/phi_scalar_simple_1.ll +++ b/polly/test/Isl/CodeGen/phi_scalar_simple_1.ll @@ -22,13 +22,13 @@ entry: br label %for.cond ; CHECK-LABEL: polly.merge_new_and_old: -; CHECK: %x.addr.0.merge = phi i32 [ %x.addr.0.final_reload, %polly.merge ], [ %x.addr.0, %for.cond ] +; CHECK: %x.addr.0.merge = phi i32 [ %x.addr.0.final_reload, %polly.merge20 ], [ %x.addr.0, %for.cond ] ; CHECK: ret i32 %x.addr.0.merge ; CHECK-LABEL: polly.start: ; CHECK-NEXT: store i32 %x, i32* %x.addr.0.phiops -; CHECK-LABEL: polly.merge: +; CHECK-LABEL: polly.merge20: ; CHECK: %x.addr.0.final_reload = load i32, i32* %x.addr.0.s2a for.cond: ; preds = %for.inc4, %entry @@ -40,30 +40,12 @@ for.cond: ; preds = %for.inc4, %entry %cmp = icmp slt i64 %indvars.iv, %tmp br i1 %cmp, label %for.body, label %for.end6 -; CHECK-LABEL: polly.stmt.for.cond{{[0-9]*}}: -; CHECK: %x.addr.0.phiops.reload[[R1:[0-9]*]] = load i32, i32* %x.addr.0.phiops -; CHECK: store i32 %x.addr.0.phiops.reload[[R1]], i32* %x.addr.0.s2a - for.body: ; preds = %for.cond ; CHECK-LABEL: polly.stmt.for.body: ; CHECK: %x.addr.0.s2a.reload[[R2:[0-9]*]] = load i32, i32* %x.addr.0.s2a ; CHECK: store i32 %x.addr.0.s2a.reload[[R2]], i32* %x.addr.1.phiops br label %for.cond1 -for.end: ; preds = %for.cond1 -; CHECK-LABEL: polly.stmt.for.end: -; CHECK-NEXT: %x.addr.1.lcssa.phiops.reload = load i32, i32* %x.addr.1.lcssa.phiops -; CHECK-NEXT: store i32 %x.addr.1.lcssa.phiops.reload, i32* %x.addr.1.lcssa.s2a[[R4:[0-9]*]] - %x.addr.1.lcssa = phi i32 [ %x.addr.1, %for.cond1 ] - br label %for.inc4 - -for.inc4: ; preds = %for.end -; CHECK-LABEL: polly.stmt.for.inc4: -; CHECK: %x.addr.1.lcssa.s2a.reload[[R5:[0-9]*]] = load i32, i32* %x.addr.1.lcssa.s2a[[R4]] -; CHECK: store i32 %x.addr.1.lcssa.s2a.reload[[R5]], i32* %x.addr.0.phiops - %indvars.iv.next = add nuw nsw i64 %indvars.iv, 1 - br label %for.cond - for.cond1: ; preds = %for.inc, %for.body ; CHECK-LABEL: polly.stmt.for.cond1: ; CHECK: %x.addr.1.phiops.reload = load i32, i32* %x.addr.1.phiops @@ -74,9 +56,6 @@ for.cond1: ; preds = %for.inc, %for.body %exitcond = icmp ne i32 %j.0, %N br i1 %exitcond, label %for.body3, label %for.end -for.body3: ; preds = %for.cond1 - br label %for.inc - for.inc: ; preds = %for.body3 ; CHECK-LABEL: polly.stmt.for.inc: ; CHECK: %x.addr.1.s2a.reload[[R3:[0-9]*]] = load i32, i32* %x.addr.1.s2a @@ -88,6 +67,23 @@ for.inc: ; preds = %for.body3 %inc = add nsw i32 %j.0, 1 br label %for.cond1 +for.end: ; preds = %for.cond1 +; CHECK-LABEL: polly.stmt.for.end: +; CHECK-NEXT: %x.addr.1.lcssa.phiops.reload = load i32, i32* %x.addr.1.lcssa.phiops +; CHECK-NEXT: store i32 %x.addr.1.lcssa.phiops.reload, i32* %x.addr.1.lcssa.s2a[[R4:[0-9]*]] + %x.addr.1.lcssa = phi i32 [ %x.addr.1, %for.cond1 ] + br label %for.inc4 + +for.inc4: ; preds = %for.end +; CHECK-LABEL: polly.stmt.for.inc4: +; CHECK: %x.addr.1.lcssa.s2a.reload[[R5:[0-9]*]] = load i32, i32* %x.addr.1.lcssa.s2a[[R4]] +; CHECK: store i32 %x.addr.1.lcssa.s2a.reload[[R5]], i32* %x.addr.0.phiops + %indvars.iv.next = add nuw nsw i64 %indvars.iv, 1 + br label %for.cond + +for.body3: ; preds = %for.cond1 + br label %for.inc + for.end6: ; preds = %for.cond ret i32 %x.addr.0 } diff --git a/polly/test/Isl/CodeGen/two-scops-in-row.ll b/polly/test/Isl/CodeGen/two-scops-in-row.ll index 25bfe6727020..fd8f7f7cae4c 100644 --- a/polly/test/Isl/CodeGen/two-scops-in-row.ll +++ b/polly/test/Isl/CodeGen/two-scops-in-row.ll @@ -5,10 +5,9 @@ target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" ; SCALAR: if (1) ; SCALAR: { -; SCALAR: for (int c0 = 0; c0 <= -Scalar0.val + 99; c0 += 1) +; SCALAR: Stmt_for_1(0); +; SCALAR: for (int c0 = 1; c0 <= -Scalar0.val + 99; c0 += 1) ; SCALAR: Stmt_for_1(c0); -; SCALAR: if (Scalar0.val >= 100) -; SCALAR: Stmt_for_1(0); ; SCALAR: } ; SCALAR: if (1) diff --git a/polly/test/ScopInfo/NonAffine/div_backedge.ll b/polly/test/ScopInfo/NonAffine/div_backedge.ll new file mode 100644 index 000000000000..e6db95903f91 --- /dev/null +++ b/polly/test/ScopInfo/NonAffine/div_backedge.ll @@ -0,0 +1,42 @@ +; RUN: opt %loadPolly -polly-detect-unprofitable -polly-scops -analyze < %s | FileCheck %s +; +; void foo(float *A) { +; for (long i = 1;; i++) { +; A[i] += 1; +; if (i / 7 == 4) +; break; +; } +; } +; +; CHECK: Domain := +; CHECK: { Stmt_for_body[i0] : i0 <= 27 and i0 >= 0 }; +; +target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" + +define void @foo(float* %A) { +entry: + br label %for.cond + +for.cond: ; preds = %for.inc, %entry + %i.0 = phi i64 [ 1, %entry ], [ %inc, %for.inc ] + br label %for.body + +for.body: ; preds = %for.cond + %arrayidx0 = getelementptr inbounds float, float* %A, i64 %i.0 + %tmp0 = load float, float* %arrayidx0, align 4 + %add0 = fadd float %tmp0, 2.000000e+00 + store float %add0, float* %arrayidx0, align 4 + %rem1 = sdiv i64 %i.0, 7 + %tobool = icmp eq i64 %rem1, 4 + br i1 %tobool, label %for.end, label %if.end + +if.end: ; preds = %for.body, %if.then + br label %for.inc + +for.inc: ; preds = %if.end + %inc = add nuw nsw i64 %i.0, 1 + br label %for.cond + +for.end: ; preds = %for.cond + ret void +} diff --git a/polly/test/ScopInfo/NonAffine/div_domain.ll b/polly/test/ScopInfo/NonAffine/div_domain.ll new file mode 100644 index 000000000000..af6c00d69454 --- /dev/null +++ b/polly/test/ScopInfo/NonAffine/div_domain.ll @@ -0,0 +1,54 @@ +; RUN: opt %loadPolly -polly-scops -analyze -polly-detect-unprofitable < %s | FileCheck %s +; +; void foo(float *A) { +; for (long i = 0; i < 16; i++) { +; A[i] += 1; +; if (i / 2 == 3) +; A[i] += 2; +; } +; } +; +; CHECK: Stmt_for_body +; CHECK: Domain := +; CHECK: { Stmt_for_body[i0] : i0 <= 15 and i0 >= 0 }; +; CHECK: Stmt_if_then +; CHECK: Domain := +; CHECK: { Stmt_if_then[i0] : i0 <= 7 and i0 >= 6 }; +; +target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" + +define void @foo(float* %A) { +entry: + br label %for.cond + +for.cond: ; preds = %for.inc, %entry + %i.0 = phi i64 [ 0, %entry ], [ %inc, %for.inc ] + %exitcond = icmp ne i64 %i.0, 16 + br i1 %exitcond, label %for.body, label %for.end + +for.body: ; preds = %for.cond + %arrayidx0 = getelementptr inbounds float, float* %A, i64 %i.0 + %tmp0 = load float, float* %arrayidx0, align 4 + %add0 = fadd float %tmp0, 2.000000e+00 + store float %add0, float* %arrayidx0, align 4 + %rem1 = sdiv i64 %i.0, 2 + %tobool = icmp ne i64 %rem1, 3 + br i1 %tobool, label %if.end, label %if.then + +if.then: ; preds = %for.body + %arrayidx = getelementptr inbounds float, float* %A, i64 %i.0 + %tmp = load float, float* %arrayidx, align 4 + %add = fadd float %tmp, 2.000000e+00 + store float %add, float* %arrayidx, align 4 + br label %if.end + +if.end: ; preds = %for.body, %if.then + br label %for.inc + +for.inc: ; preds = %if.end + %inc = add nuw nsw i64 %i.0, 1 + br label %for.cond + +for.end: ; preds = %for.cond + ret void +} diff --git a/polly/test/ScopInfo/NonAffine/modulo_backedge.ll b/polly/test/ScopInfo/NonAffine/modulo_backedge.ll new file mode 100644 index 000000000000..3d35c2105666 --- /dev/null +++ b/polly/test/ScopInfo/NonAffine/modulo_backedge.ll @@ -0,0 +1,44 @@ +; RUN: opt %loadPolly -polly-scops -polly-detect-unprofitable -analyze < %s | FileCheck %s +; +; TODO: The new domain generation cannot handle modulo domain constraints, +; hence modulo handling has been disabled completely. Once this is +; resolved this test should work again. +; CHECK-NOT: Access +; +; void foo(float *A) { +; for (long i = 1;; i++) { +; A[i] += 1; +; if (i % 7 == 0) +; break; +; } +; } +; +target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" + +define void @foo(float* %A) { +entry: + br label %for.cond + +for.cond: ; preds = %for.inc, %entry + %i.0 = phi i64 [ 1, %entry ], [ %inc, %for.inc ] + br label %for.body + +for.body: ; preds = %for.cond + %arrayidx0 = getelementptr inbounds float, float* %A, i64 %i.0 + %tmp0 = load float, float* %arrayidx0, align 4 + %add0 = fadd float %tmp0, 2.000000e+00 + store float %add0, float* %arrayidx0, align 4 + %rem1 = srem i64 %i.0, 7 + %tobool = icmp eq i64 %rem1, 0 + br i1 %tobool, label %for.end, label %if.end + +if.end: ; preds = %for.body, %if.then + br label %for.inc + +for.inc: ; preds = %if.end + %inc = add nuw nsw i64 %i.0, 1 + br label %for.cond + +for.end: ; preds = %for.cond + ret void +} diff --git a/polly/test/ScopInfo/NonAffine/modulo_domain.ll b/polly/test/ScopInfo/NonAffine/modulo_domain.ll new file mode 100644 index 000000000000..6b3b383d0182 --- /dev/null +++ b/polly/test/ScopInfo/NonAffine/modulo_domain.ll @@ -0,0 +1,56 @@ +; RUN: opt %loadPolly -polly-scops -polly-detect-unprofitable -analyze < %s | FileCheck %s +; +; TODO: The new domain generation cannot handle modulo domain constraints, +; hence modulo handling has been disabled completely. Once this is +; resolved this test should work again. Until then we approximate the +; whole loop body. +; CHECK: Stmt_for_body__TO__if_end +; CHECK: Domain := +; CHECK: { Stmt_for_body__TO__if_end[i0] : i0 <= 15 and i0 >= 0 }; + +; +; void foo(float *A) { +; for (long i = 0; i < 16; i++) { +; A[i] += 1; +; if (i % 2) +; A[i] += 2; +; } +; } +; +target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" + +define void @foo(float* %A) { +entry: + br label %for.cond + +for.cond: ; preds = %for.inc, %entry + %i.0 = phi i64 [ 0, %entry ], [ %inc, %for.inc ] + %exitcond = icmp ne i64 %i.0, 16 + br i1 %exitcond, label %for.body, label %for.end + +for.body: ; preds = %for.cond + %arrayidx0 = getelementptr inbounds float, float* %A, i64 %i.0 + %tmp0 = load float, float* %arrayidx0, align 4 + %add0 = fadd float %tmp0, 2.000000e+00 + store float %add0, float* %arrayidx0, align 4 + %rem1 = srem i64 %i.0, 2 + %tobool = icmp eq i64 %rem1, 0 + br i1 %tobool, label %if.end, label %if.then + +if.then: ; preds = %for.body + %arrayidx = getelementptr inbounds float, float* %A, i64 %i.0 + %tmp = load float, float* %arrayidx, align 4 + %add = fadd float %tmp, 2.000000e+00 + store float %add, float* %arrayidx, align 4 + br label %if.end + +if.end: ; preds = %for.body, %if.then + br label %for.inc + +for.inc: ; preds = %if.end + %inc = add nuw nsw i64 %i.0, 1 + br label %for.cond + +for.end: ; preds = %for.cond + ret void +} diff --git a/polly/test/ScopInfo/NonAffine/non-affine-loop-condition-dependent-access_2.ll b/polly/test/ScopInfo/NonAffine/non-affine-loop-condition-dependent-access_2.ll index 137c54c423a5..6ba4db4ca195 100644 --- a/polly/test/ScopInfo/NonAffine/non-affine-loop-condition-dependent-access_2.ll +++ b/polly/test/ScopInfo/NonAffine/non-affine-loop-condition-dependent-access_2.ll @@ -19,35 +19,34 @@ ; INNERMOST: Region: %bb15---%bb13 ; INNERMOST: Max Loop Depth: 1 ; INNERMOST: p0: {0,+,{0,+,1}<%bb11>}<%bb13> -; INNERMOST: p1: {0,+,{0,+,-1}<%bb11>}<%bb13> -; INNERMOST: p2: {0,+,4}<%bb11> -; INNERMOST: p3: {0,+,4}<%bb13> -; INNERMOST: p4: {0,+,{0,+,4}<%bb11>}<%bb13> +; INNERMOST: p1: {0,+,4}<%bb11> +; INNERMOST: p2: {0,+,4}<%bb13> +; INNERMOST: p3: {0,+,{0,+,4}<%bb11>}<%bb13> ; INNERMOST: Alias Groups (0): ; INNERMOST: n/a ; INNERMOST: Statements { ; INNERMOST: Stmt_bb16 ; INNERMOST: Domain := -; TODO-INNERMOST: [p_0, p_1, p_2, p_3, p_4] -> { Stmt_bb16[i0] : i0 <= 1023 - p_0 and i0 >= 0 } +; INNERMOST: [p_0, p_1, p_2, p_3] -> { Stmt_bb16[i0] : i0 <= 1023 - p_0 and i0 >= 0 } ; INNERMOST: Schedule := -; TODO-INNERMOST: [p_0, p_1, p_2, p_3, p_4] -> { Stmt_bb16[i0] -> [i0] } +; INNERMOST: [p_0, p_1, p_2, p_3] -> { Stmt_bb16[i0] -> [0, i0] } ; INNERMOST: ReadAccess := [Reduction Type: NONE] [Scalar: 0] -; INNERMOST: [p_0, p_1, p_2, p_3, p_4] -> { Stmt_bb16[i0] -> MemRef_A[o0] : 4o0 = p_2 }; +; INNERMOST: [p_0, p_1, p_2, p_3] -> { Stmt_bb16[i0] -> MemRef_A[o0] : 4o0 = p_1 }; ; INNERMOST: ReadAccess := [Reduction Type: NONE] [Scalar: 0] -; INNERMOST: [p_0, p_1, p_2, p_3, p_4] -> { Stmt_bb16[i0] -> MemRef_A[o0] : 4o0 = p_3 }; +; INNERMOST: [p_0, p_1, p_2, p_3] -> { Stmt_bb16[i0] -> MemRef_A[o0] : 4o0 = p_2 }; ; INNERMOST: ReadAccess := [Reduction Type: +] [Scalar: 0] -; INNERMOST: [p_0, p_1, p_2, p_3, p_4] -> { Stmt_bb16[i0] -> MemRef_A[o0] : 4o0 = p_4 + 4i0 }; +; INNERMOST: [p_0, p_1, p_2, p_3] -> { Stmt_bb16[i0] -> MemRef_A[o0] : 4o0 = p_3 + 4i0 }; ; INNERMOST: MustWriteAccess := [Reduction Type: +] [Scalar: 0] -; INNERMOST: [p_0, p_1, p_2, p_3, p_4] -> { Stmt_bb16[i0] -> MemRef_A[o0] : 4o0 = p_4 + 4i0 }; +; INNERMOST: [p_0, p_1, p_2, p_3] -> { Stmt_bb16[i0] -> MemRef_A[o0] : 4o0 = p_3 + 4i0 }; ; INNERMOST: Stmt_bb26 ; INNERMOST: Domain := -; INNERMOST: [p_0, p_1, p_2, p_3, p_4] -> { Stmt_bb26[] }; +; INNERMOST: [p_0, p_1, p_2, p_3] -> { Stmt_bb26[] : p_0 <= 1024 }; ; INNERMOST: Schedule := -; INNERMOST: [p_0, p_1, p_2, p_3, p_4] -> { Stmt_bb26[] -> [1, 0] }; +; INNERMOST: [p_0, p_1, p_2, p_3] -> { Stmt_bb26[] -> [1, 0] }; ; INNERMOST: MustWriteAccess := [Reduction Type: NONE] [Scalar: 1] -; INNERMOST: [p_0, p_1, p_2, p_3, p_4] -> { Stmt_bb26[] -> MemRef_indvars_iv_next6[] }; +; INNERMOST: [p_0, p_1, p_2, p_3] -> { Stmt_bb26[] -> MemRef_indvars_iv_next6[] }; ; INNERMOST: MustWriteAccess := [Reduction Type: NONE] [Scalar: 1] -; INNERMOST: [p_0, p_1, p_2, p_3, p_4] -> { Stmt_bb26[] -> MemRef_indvars_iv_next4[] }; +; INNERMOST: [p_0, p_1, p_2, p_3] -> { Stmt_bb26[] -> MemRef_indvars_iv_next4[] }; ; INNERMOST: } ; ; ALL: Function: f diff --git a/polly/test/ScopInfo/NonAffine/non-affine-loop-condition-dependent-access_3.ll b/polly/test/ScopInfo/NonAffine/non-affine-loop-condition-dependent-access_3.ll index 222e28c75aee..5047a80c668f 100644 --- a/polly/test/ScopInfo/NonAffine/non-affine-loop-condition-dependent-access_3.ll +++ b/polly/test/ScopInfo/NonAffine/non-affine-loop-condition-dependent-access_3.ll @@ -46,7 +46,7 @@ ; INNERMOST: [p_0, p_1, p_2] -> { Stmt_bb16[i0] -> MemRef_A[i0] }; ; INNERMOST: Stmt_bb26 ; INNERMOST: Domain := -; INNERMOST: [p_0, p_1, p_2] -> { Stmt_bb26[] }; +; INNERMOST: [p_0, p_1, p_2] -> { Stmt_bb26[] : p_0 >= 0 }; ; INNERMOST: Schedule := ; INNERMOST: [p_0, p_1, p_2] -> { Stmt_bb26[] -> [1, 0] }; ; INNERMOST: MustWriteAccess := [Reduction Type: NONE] [Scalar: 1] diff --git a/polly/test/ScopInfo/isl_trip_count_01.ll b/polly/test/ScopInfo/isl_trip_count_01.ll index 1388dfb7314a..60ac38cfb501 100644 --- a/polly/test/ScopInfo/isl_trip_count_01.ll +++ b/polly/test/ScopInfo/isl_trip_count_01.ll @@ -1,6 +1,6 @@ ; RUN: opt %loadPolly -polly-detect-unprofitable -polly-allow-non-scev-backedge-taken-count -polly-scops -analyze < %s | FileCheck %s ; -; CHECK: [M, N] -> { Stmt_while_body[i0] : i0 >= 0 and 4i0 <= -M + N } +; CHECK: [M, N] -> { Stmt_while_body[i0] : i0 >= 1 and 4i0 <= -M + N; Stmt_while_body[0] } ; ; void f(int *A, int N, int M) { ; int i = 0; diff --git a/polly/test/ScopInfo/loop_affine_bound_0.ll b/polly/test/ScopInfo/loop_affine_bound_0.ll index 1e409c1a2103..52269dc82b44 100644 --- a/polly/test/ScopInfo/loop_affine_bound_0.ll +++ b/polly/test/ScopInfo/loop_affine_bound_0.ll @@ -63,8 +63,6 @@ return: ; preds = %bb.nph8, %bb3, %ent ; CHECK-DAG: i1 >= 0 ; CHECK-DAG: and ; CHECK-DAG: i1 <= 1 + 5N -; CHECK-DAG: and -; CHECK-DAG: N >= 0 ; CHECK: } ; CHECK: Schedule := ; CHECK: [N, M] -> { Stmt_bb1[i0, i1] -> [i0, i1] }; diff --git a/polly/test/ScopInfo/loop_affine_bound_1.ll b/polly/test/ScopInfo/loop_affine_bound_1.ll index 0da94f79e62a..5d137b67f7bf 100644 --- a/polly/test/ScopInfo/loop_affine_bound_1.ll +++ b/polly/test/ScopInfo/loop_affine_bound_1.ll @@ -16,7 +16,7 @@ entry: %2 = or i64 %0, 3 ; [#uses=1] %3 = add nsw i64 %2, %1 ; [#uses=1] %4 = icmp sgt i64 %3, 0 ; [#uses=1] - br i1 %4, label %bb.nph8, label %return + br i1 true, label %bb.nph8, label %return bb1: ; preds = %bb2.preheader, %bb1 %indvar = phi i64 [ 0, %bb2.preheader ], [ %indvar.next, %bb1 ] ; [#uses=2] @@ -55,18 +55,18 @@ return: ; preds = %bb3, %entry ; CHECK: Stmt_bb1 ; CHECK: Domain := ; CHECK: [N, M] -> { Stmt_bb1[i0, i1] : -; CHECK-DAG: i0 >= 0 +; CHECK-DAG: i0 >= 1 ; CHECK-DAG: and ; CHECK-DAG: i0 <= 2 + 4N + 7M ; CHECK-DAG: and ; CHECK-DAG: i1 >= 0 ; CHECK-DAG: and ; CHECK-DAG: i1 <= 1 + 5N - i0 +; CHECK-DAG: Stmt_bb1[0, i1] : +; CHECK-DAG: i1 >= 0 ; CHECK-DAG: and -; CHECK-DAG: i0 <= 1 + 5N +; CHECK-DAG: i1 <= 1 + 5N ; CHECK: } -; CHECK: Schedule := -; CHECK: [N, M] -> { Stmt_bb1[i0, i1] -> [i0, i1] }; ; CHECK: MustWriteAccess := [Reduction Type: NONE] ; CHECK: [N, M] -> { Stmt_bb1[i0, i1] -> MemRef_a[129i0 + 128i1] }; ; CHECK: } diff --git a/polly/test/ScopInfo/loop_affine_bound_2.ll b/polly/test/ScopInfo/loop_affine_bound_2.ll index 41f67138d1b9..3f6835ec9504 100644 --- a/polly/test/ScopInfo/loop_affine_bound_2.ll +++ b/polly/test/ScopInfo/loop_affine_bound_2.ll @@ -73,8 +73,6 @@ return: ; preds = %bb3, %entry ; CHECK-DAG: i1 >= 0 ; CHECK-DAG: and ; CHECK-DAG: i1 <= 10 + 5N - 6M - 4i0 -; CHECK-DAG: and -; CHECK-DAG: 4i0 <= 10 + 5N - 6M ; CHECK-DAG: } ; CHECK: Schedule := ; CHECK: [N, M] -> { Stmt_bb1[i0, i1] -> [i0, i1] diff --git a/polly/test/ScopInfo/non_affine_region_1.ll b/polly/test/ScopInfo/non_affine_region_1.ll index 0af44cd73517..895851071452 100644 --- a/polly/test/ScopInfo/non_affine_region_1.ll +++ b/polly/test/ScopInfo/non_affine_region_1.ll @@ -36,8 +36,10 @@ ; CHECK-DAG: i0 <= 1023 ; CHECK: }; ; CHECK-NEXT: Schedule := -; CHECK-NEXT: [b] -> { Stmt_bb10__TO__bb18[i0] -> [i0, 3] }; -; CHECK-NEXT: ReadAccess := [Reduction Type: NONE] [Scalar: 1] +; TODO: We build a complicated representation of the domain that will also complicate the schedule. +; Once the domain is as simple as shown above this test should fail and this TODO can be removed. +; CHECK-NOT: [b] -> { Stmt_bb10__TO__bb18[i0] -> [i0, 3] }; +; CHECK: ReadAccess := [Reduction Type: NONE] [Scalar: 1] ; CHECK-NEXT: [b] -> { Stmt_bb10__TO__bb18[i0] -> MemRef_x_1__phi[] } ; CHECK-NOT: [Scalar: 1] ; diff --git a/polly/test/ScopInfo/pointer-type-expressions.ll b/polly/test/ScopInfo/pointer-type-expressions.ll index 44c7de37fda0..df8c3d491803 100644 --- a/polly/test/ScopInfo/pointer-type-expressions.ll +++ b/polly/test/ScopInfo/pointer-type-expressions.ll @@ -43,9 +43,9 @@ return: ; CHECK: Stmt_store ; CHECK: Domain := ; CHECK: [P, N] -> { Stmt_store[i0] : -; CHECK: (P <= -1 and i0 >= 0 and i0 <= -1 + N) -; CHECK: or -; CHECK: (P >= 1 and i0 >= 0 and i0 <= -1 + N) +; CHECK-DAG: (P <= -1 and i0 >= 0 and i0 <= -1 + N) +; CHECK-DAG: or +; CHECK-DAG: (P >= 1 and i0 >= 0 and i0 <= -1 + N) ; CHECK: }; ; CHECK: Schedule := ; CHECK: [P, N] -> { Stmt_store[i0] -> [i0] : P <= -1 or P >= 1 }; diff --git a/polly/test/ScopInfo/redundant_parameter_constraint.ll b/polly/test/ScopInfo/redundant_parameter_constraint.ll new file mode 100644 index 000000000000..5da19c0fc1ad --- /dev/null +++ b/polly/test/ScopInfo/redundant_parameter_constraint.ll @@ -0,0 +1,43 @@ +; RUN: opt %loadPolly -polly-detect-unprofitable -polly-scops -analyze < %s | FileCheck %s +; +; The constraint that r2 has to be bigger than r1 is implicitly containted in +; the domain, hence we do not want to see it explicitly. +; +; CHECK-NOT: r2 >= 1 + r1 +; +; void wraps(int *A, int p, short q, char r1, char r2) { +; for (char i = r1; i < r2; i++) +; A[p + q] = A[(int)r1 + (int)r2]; +; } +; +target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" + +define void @wraps(i32* %A, i32 %p, i16 signext %q, i8 signext %r1, i8 signext %r2) { +entry: + br label %for.cond + +for.cond: ; preds = %for.inc, %entry + %i.0 = phi i8 [ %r1, %entry ], [ %inc, %for.inc ] + %cmp = icmp slt i8 %i.0, %r2 + br i1 %cmp, label %for.body, label %for.end + +for.body: ; preds = %for.cond + %conv3 = sext i8 %r1 to i64 + %conv4 = sext i8 %r2 to i64 + %add = add nsw i64 %conv3, %conv4 + %arrayidx = getelementptr inbounds i32, i32* %A, i64 %add + %tmp = load i32, i32* %arrayidx, align 4 + %conv5 = sext i16 %q to i32 + %add6 = add nsw i32 %conv5, %p + %idxprom7 = sext i32 %add6 to i64 + %arrayidx8 = getelementptr inbounds i32, i32* %A, i64 %idxprom7 + store i32 %tmp, i32* %arrayidx8, align 4 + br label %for.inc + +for.inc: ; preds = %for.body + %inc = add i8 %i.0, 1 + br label %for.cond + +for.end: ; preds = %for.cond + ret void +} diff --git a/polly/test/ScopInfo/two-loops-right-after-each-other.ll b/polly/test/ScopInfo/two-loops-right-after-each-other.ll index 5230c3581883..5c329e067cb9 100644 --- a/polly/test/ScopInfo/two-loops-right-after-each-other.ll +++ b/polly/test/ScopInfo/two-loops-right-after-each-other.ll @@ -3,10 +3,22 @@ ; CHECK: Stmt_loop_1 ; CHECK: Domain := -; CHECK: [N] -> { Stmt_loop_1[i0] : i0 >= 0 and i0 <= 101 and N <= 100 }; +; CHECK: [N] -> { Stmt_loop_1[i0] : +; CHECK-DAG: i0 >= 0 +; CHECK-DAG: and +; CHECK-DAG: i0 <= 101 +; CHECK-DAG: and +; CHECK-DAG: N <= 100 +; CHECK-DAG: }; ; CHECK: Stmt_loop_2 ; CHECK: Domain := -; CHECK: [N] -> { Stmt_loop_2[i0] : i0 >= 0 and i0 <= 301 and N <= 100 }; +; CHECK: [N] -> { Stmt_loop_2[i0] : +; CHECK-DAG: i0 >= 0 +; CHECK-DAG: and +; CHECK-DAG: i0 <= 301 +; CHECK-DAG: and +; CHECK-DAG: N <= 100 +; CHECK-DAG: }; define void @foo(float* %A, i64 %N) { entry: diff --git a/polly/test/ScopInfo/unsigned-condition.ll b/polly/test/ScopInfo/unsigned-condition.ll index eadffb4d3e4a..d9ee1897bfbc 100644 --- a/polly/test/ScopInfo/unsigned-condition.ll +++ b/polly/test/ScopInfo/unsigned-condition.ll @@ -38,7 +38,11 @@ return: ; CHECK: Stmt_store ; CHECK: Domain := ; CHECK: [P, N] -> { Stmt_store[i0] : -; CHECK: i0 >= 0 and i0 <= -1 + N and P >= 42 +; CHECK-DAG: i0 >= 0 +; CHECK-DAG: and +; CHECK-DAG: i0 <= -1 + N +; CHECK-DAG: and +; CHECK-DAG: P >= 42 ; CHECK: }; ; CHECK: Schedule := ; CHECK: [P, N] -> { Stmt_store[i0] -> [i0] }; -- GitLab