Index: include/polly/ScopInfo.h =================================================================== --- include/polly/ScopInfo.h +++ include/polly/ScopInfo.h @@ -554,9 +554,6 @@ /// 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. @@ -910,6 +907,14 @@ AliasAnalysis &AA, DominatorTree &DT, isl_ctx *ctx); + /// @brief Add loop carried constraints to the domain of each block. + /// + /// @param LI The LoopInfo analysis to argue about 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 addLoopConstraintsToDomains(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. @@ -919,6 +924,15 @@ 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. Index: lib/Analysis/ScopInfo.cpp =================================================================== --- lib/Analysis/ScopInfo.cpp +++ lib/Analysis/ScopInfo.cpp @@ -31,6 +31,7 @@ #include "llvm/ADT/StringExtras.h" #include "llvm/ADT/PostOrderIterator.h" #include "llvm/Analysis/AliasAnalysis.h" +#include "llvm/Analysis/LoopIterator.h" #include "llvm/Analysis/LoopInfo.h" #include "llvm/Analysis/RegionIterator.h" #include "llvm/Analysis/ScalarEvolutionExpressions.h" @@ -812,117 +813,63 @@ 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); +/// @brief Add @p bset to the set @p User if @p bset is unbounded. +static isl_stat collectUnboundedParts(__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; } -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. - addLoopTripCountToDomain(L); - isl_pw_aff_free(IV); - } +/// @brief Return @p S without basic sets that are unbounded in dimension @p +/// Dim. +static __isl_give std::pair +isolateBoundedParts(__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 @p Dim as we are not interested in + // them. + OnlyDimS = + isl_set_remove_dims(OnlyDimS, isl_dim_set, Dim + 1, NumDimsS - Dim - 1); + + // Create upper bounds for dimensions smaller than @p Dim as we are not + // interested in + // them. + OnlyDimS = isl_set_insert_dims(OnlyDimS, isl_dim_param, 0, Dim); + for (unsigned u = 0, e = Dim; u < e; 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); } - isl_local_space_free(LocalSpace); + isl_set *BoundedParts = isl_set_empty(isl_set_get_space(OnlyDimS)); + isl_set_foreach_basic_set(OnlyDimS, collectUnboundedParts, &BoundedParts); + isl_set_free(OnlyDimS); + + BoundedParts = isl_set_insert_dims(BoundedParts, isl_dim_set, Dim + 1, + NumDimsS - Dim - 1); + BoundedParts = isl_set_remove_dims(BoundedParts, isl_dim_param, 0, Dim); + + return std::make_pair(isl_set_subtract(S, isl_set_copy(BoundedParts)), + BoundedParts); } 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); } @@ -1471,20 +1418,32 @@ return BI->getSuccessor(idx); } +static inline Loop *getRegionLoop(Region *R, LoopInfo &LI) { + Loop *L = LI.getLoopFor(R->getEntry()); + while (L && R->contains(L)) + L = L->getParentLoop(); + return L; +} + /// @brief Return the smallest loop surrounding @p RN. static inline Loop *getRegionNodeLoop(RegionNode *RN, LoopInfo &LI) { if (!RN->isSubRegion()) return LI.getLoopFor(RN->getNodeAs()); Region *NonAffineSubRegion = RN->getNodeAs(); - Loop *L = LI.getLoopFor(NonAffineSubRegion->getEntry()); - while (L && NonAffineSubRegion->contains(L)) - L = L->getParentLoop(); - return L; + return getRegionLoop(NonAffineSubRegion, LI); } ///} +static __isl_give isl_set * +createFirstIterationDomain(__isl_take isl_space *Space, int Dim) { + auto *Domain = isl_set_universe(Space); + if (Dim >= 0) + Domain = isl_set_fix_si(Domain, isl_dim_set, Dim, 0); + return Domain; +} + isl_set *Scop::getDomainConditions(ScopStmt *Stmt) { BasicBlock *BB = Stmt->isBlockStmt() ? Stmt->getBasicBlock() : Stmt->getRegion()->getEntry(); @@ -1500,6 +1459,8 @@ DomainMap[EntryBB] = S; buildDomainsWithBranchConstraints(R, LI, SD, DT); + addLoopConstraintsToDomains(LI, SD, DT); + propagateDomainConstraints(R, LI, SD, DT); } void Scop::buildDomainsWithBranchConstraints(Region *R, LoopInfo &LI, @@ -1605,6 +1566,205 @@ } } +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) { + + const ScopDetection::BoxedLoopsSetTy *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]; + DEBUG(dbgs() << "\tVisit: " << BB->getName() << " : "; + isl_set_dump(Domain)); + 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; + if (!getRegion().contains(PredBB)) + PredBBDom = isl_set_universe(isl_set_get_space(PredDom)); + + if (!PredBBDom) { + 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_remove_dims(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); + } + + 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, hence: +/// [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 and its 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 all 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()); + } +} + +void Scop::addLoopConstraintsToDomains(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. Later we propagate them through the whole loop. + + 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); + isl_set_free(ConditionSets[1 - idx]); + + BackedgeCondition = ConditionSets[idx]; + } + + auto *&HeaderBBDom = DomainMap[HeaderBB]; + isl_set *FirstIteration = + createFirstIterationDomain(isl_set_get_space(HeaderBBDom), LoopDepth); + + // Handle pseudo backedges that are never taken. + if (isl_set_is_empty(BackedgeCondition)) { + isl_set_free(BackedgeCondition); + BackedgeCondition = FirstIteration; + } else { + isl_space *SetSpace = isl_set_get_space(HeaderBBDom); + isl_map *NextIterationMap = createNextIterationMap(SetSpace, LoopDepth); + + int LatchLoopDepth = getRelativeLoopDepth(LI.getLoopFor(LatchBB)); + assert(LatchLoopDepth >= LoopDepth); + BackedgeCondition = + isl_set_remove_dims(BackedgeCondition, isl_dim_set, LoopDepth + 1, + LatchLoopDepth - LoopDepth); + + auto Parts = isolateBoundedParts(BackedgeCondition, LoopDepth); + + 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; Index: test/Isl/CodeGen/OpenMP/loop-body-references-outer-values-3.ll =================================================================== --- test/Isl/CodeGen/OpenMP/loop-body-references-outer-values-3.ll +++ 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 @@ ; 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 Index: test/Isl/CodeGen/OpenMP/reference-preceeding-loop.ll =================================================================== --- test/Isl/CodeGen/OpenMP/reference-preceeding-loop.ll +++ 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" Index: test/Isl/CodeGen/two-scops-in-row.ll =================================================================== --- test/Isl/CodeGen/two-scops-in-row.ll +++ test/Isl/CodeGen/two-scops-in-row.ll @@ -5,10 +5,9 @@ ; 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) Index: test/ScopInfo/NonAffine/non-affine-loop-condition-dependent-access_2.ll =================================================================== --- test/ScopInfo/NonAffine/non-affine-loop-condition-dependent-access_2.ll +++ test/ScopInfo/NonAffine/non-affine-loop-condition-dependent-access_2.ll @@ -18,27 +18,33 @@ ; INNERMOST: Function: f ; INNERMOST: Region: %bb15---%bb26 ; INNERMOST: Max Loop Depth: 1 +;OLDINNERMOST: p0: {0,+,{0,+,1}<%bb11>}<%bb13> +;OLDINNERMOST: p1: {0,+,{0,+,-1}<%bb11>}<%bb13> +;OLDINNERMOST: p2: {0,+,4}<%bb11> +;OLDINNERMOST: p3: {0,+,4}<%bb13> +;OLDINNERMOST: p4: {0,+,{0,+,4}<%bb11>}<%bb13> ; 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 } +;OLDINNERMOST: [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] } +;OLDINNERMOST: [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] -> [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: } ; ; ALL: Function: f Index: test/ScopInfo/isl_trip_count_01.ll =================================================================== --- test/ScopInfo/isl_trip_count_01.ll +++ 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; Index: test/ScopInfo/isl_trip_count_02.ll =================================================================== --- test/ScopInfo/isl_trip_count_02.ll +++ test/ScopInfo/isl_trip_count_02.ll @@ -1,6 +1,8 @@ ; RUN: opt %loadPolly -polly-detect-unprofitable -polly-allow-non-scev-backedge-taken-count -polly-scops -analyze < %s | FileCheck %s ; -; CHECK: [M, N] -> { Stmt_for_body[i0] : i0 >= 0 and N <= -1 + M }; +; TODO: This test is currently broken as it has an unconditionally infinite loop. +; +; CHECK-NOT: Stmt_for_body ; ; void f(int *A, int N, int M) { ; for (int i = M; i > N; i++) Index: test/ScopInfo/loop_affine_bound_0.ll =================================================================== --- test/ScopInfo/loop_affine_bound_0.ll +++ test/ScopInfo/loop_affine_bound_0.ll @@ -63,8 +63,6 @@ ; 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] }; Index: test/ScopInfo/loop_affine_bound_1.ll =================================================================== --- test/ScopInfo/loop_affine_bound_1.ll +++ test/ScopInfo/loop_affine_bound_1.ll @@ -16,7 +16,7 @@ %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 @@ ; 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: } Index: test/ScopInfo/loop_affine_bound_2.ll =================================================================== --- test/ScopInfo/loop_affine_bound_2.ll +++ test/ScopInfo/loop_affine_bound_2.ll @@ -73,8 +73,6 @@ ; 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] Index: test/ScopInfo/pointer-type-expressions.ll =================================================================== --- test/ScopInfo/pointer-type-expressions.ll +++ test/ScopInfo/pointer-type-expressions.ll @@ -38,9 +38,9 @@ ; 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 }; Index: test/ScopInfo/redundant_parameter_constraint.ll =================================================================== --- /dev/null +++ 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 +} Index: test/ScopInfo/unsigned-condition.ll =================================================================== --- test/ScopInfo/unsigned-condition.ll +++ test/ScopInfo/unsigned-condition.ll @@ -38,7 +38,11 @@ ; 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] };