Replace ScalarEvolution based domain generation

[jdoerfert]

Summary

This patch replaces the second legacy part of the domain generation, namely
the ScalarEvolution part that was used to obtain loop bounds. Alternatively,
we 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.

                                                                                                                                             
- 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.

[jdoerfert]

Some comments to understand the test changes.

test/Isl/CodeGen/OpenMP/loop-body-references-outer-values-3.ll

12 ↗

(On Diff #33613)

There was an if (cols >= 2) around this loop and a Stmt_for_body(0) in the else case.
The conditional is gone in the new version.

42 ↗

(On Diff #33613)

There was an if (cols >= 2) around this loop and a Stmt_for_body(0) in the else case.
The conditional is gone in the new version.

test/Isl/CodeGen/OpenMP/reference-preceeding-loop.ll

17 ↗

(On Diff #33613)

This test was funny, p_0 is actually -p_2, hence the loop had one iteration only, namely Stmt_while_body(0)

test/Isl/CodeGen/two-scops-in-row.ll

11 ↗

(On Diff #33613)

Similar to the case above, Stmt_for_1(0) is always executed here but it is explicit in the new version

test/ScopInfo/isl_trip_count_01.ll

3 ↗

(On Diff #33613)

Here the original test case is actually broken. If you look at the code but ignore the entry node (it is not in the SCoP!) than you see that for N=0 and M=1 we would execute while_body but the domain says otherwise.

test/ScopInfo/loop_affine_bound_0.ll

67 ↗

(On Diff #33613)

This is redundant as it can be extracted from the rest.

test/ScopInfo/loop_affine_bound_1.ll

19 ↗

(On Diff #33613)

This was just to make clear that the original code did not use the entry condition but just generates the wrong domain (similar to the isl_trip_count_01.ll above).

66 ↗

(On Diff #33613)

Again a broken test for M=-1 and N=0.

test/ScopInfo/loop_affine_bound_2.ll

77 ↗

(On Diff #33613)

Again redundant.

[jdoerfert]

LNT performance (with 9 tests missing) over current master

[grosser]

Hi Johannes,

this looks good for me so far.

include/polly/ScopInfo.h

559 ↗

(On Diff #33613)

It seems you could drop some parts of TempScopInfo now, as they now become unused/dead-code.

912 ↗

(On Diff #33613)

argue? This verb does not seem to make sense in this context.

lib/Analysis/ScopInfo.cpp

817 ↗

(On Diff #33613)

bset, maybe uppercase?

828 ↗

(On Diff #33613)

Please comment the contents of the return value.

831 ↗

(On Diff #33613)

Uppercase iterators?

841 ↗

(On Diff #33613)

Just removing dimensions is always fishy as this is just an operation of the internal set representation, so bad things can happen. I am not yet sure why this
is needed at all, but if it indeed is maybe you wanna use projection.

845 ↗

(On Diff #33613)

Unnecessary line break

857 ↗

(On Diff #33613)

Maybe extract out a function that collects the bounded parts of an isl_set.

1705 ↗

(On Diff #33613)

Please clarify that this propagation is not done later in this function itself, but is expected to be done afterwards (by another funciton) after this function has finished.

Also, the name of this function is misleading. Maybe call it addLoopBoundsToHeaderDomains

1729 ↗

(On Diff #33613)

Why an empty line?

1753 ↗

(On Diff #33613)

Please use project instead of remove_dims, as this does not depend on the representation of the isl_set.

[jdoerfert]

Both lnt + unit tests work, however pointer & modulo expressions are disabled

[grosser]

Hi Johannes,

some smaller comments only:

lib/Analysis/ScopDetection.cpp

351 ↗

(On Diff #34326)

Can you please commit this separately.

lib/Analysis/ScopInfo.cpp

1 ↗

(On Diff #34326)

You drop the header comment here accidentally.

lib/Support/SCEVValidator.cpp

357 ↗

(On Diff #34326)

Interesting. I do not care so much about loop bounds/conditions, but it would be sad to loose support even in array index expressions. To my understanding, srem instructions in index expressions should not cause troubles in our domain generation algorithm, no?

[jdoerfert]

lib/Support/SCEVValidator.cpp

357 ↗

(On Diff #34326)

Correct. Any ideas on how to allow those but forbid srem in domain related expressions?

[jdoerfert]

Modified according to comments.

[Meinersbur]

lib/Analysis/ScopInfo.cpp

805 ↗

(On Diff #34332)

artificial

1451 ↗

(On Diff #34332)

Function extraction into separate commit?

1546 ↗

(On Diff #34332)

Is dropping DEBUG()s this part of the patch?

1590 ↗

(On Diff #34332)

Empty comment line/missing annotations

1690 ↗

(On Diff #34332)

hence -> e.g.

1775 ↗

(On Diff #34332)

isl_set *&
LLVM coding standards: "Use auto if and only if it makes the code more readable or easier to maintain"
In any case, if you keep auto, you can also let it derive that it's a pointer.

lib/Support/SCEVValidator.cpp

356 ↗

(On Diff #34332)

Adding a boolean parameter to SCEVValidator's constructor to discriminate the two cases?

I don't get why it's not compatible for domains.

357 ↗

(On Diff #34332)

Don't commit out-commented source
Also, use

#if 0
#endif

for such cases

test/Isl/CodeGen/OpenMP/loop-body-references-outer-values-3.ll

42 ↗

(On Diff #34332)

nice

[Meinersbur]

lib/Support/SCEVValidator.cpp

356 ↗

(On Diff #34332)

We could generate a condition inside the body

for (int i = 0; i < n; i+=1) {
   if (i % 3 <= 1)
     continue;
  // stmt(i)
}

for a domain s.a. { [i] : 0 <= i < n, i - floord(i,3) * 3 <= 1 }

Alternatively, map from sequential loop counter to domain index:

for (int i_seq = 0; i_seq < (n*2/3); i_seq+=1) {
   int i = i_seq / 2 * 3 + i_seq % 2;
  // stmt(i)
}

(probably some edge-case errors in this example)

[jdoerfert]

Addressed comments by Michael

polly/trunk/include/polly/ScopInfo.h

private:
   isl_ast_build *Build;
 
   std::vector<Loop *> NestLoops;
 
   std::string BaseName;
 
   /// 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.
   ///
   /// @param tempScop       The template SCoP.
   /// @param Block          The basic block for which accesses should be
   ///                       created.
   /// @param isApproximated Flag to indicate blocks that might not be executed,
private:
 
   /// Create the static control part with a region, max loop depth of this
   /// region and parameters used in this region.
   static Scop *createFromTempScop(TempScop &TempScop, LoopInfo &LI,
                                   ScalarEvolution &SE, ScopDetection &SD,
                                   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.
   /// @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 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.
   /// @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 buildDomains(Region *R, LoopInfo &LI, ScopDetection &SD,
                     DominatorTree &DT);

polly/trunk/lib/Analysis/ScopInfo.cpp

 
 void ScopStmt::realignParams() {
   for (MemoryAccess *MA : *this)
     MA->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<isl_set **>(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) {
   case ICmpInst::ICMP_EQ:
     return isl_pw_aff_eq_set(L, R);
   case ICmpInst::ICMP_NE:
     return isl_pw_aff_ne_set(L, R);
   case ICmpInst::ICMP_SLT:
isl_set *AlternativeCondSet =
       isl_set_complement(isl_set_copy(ConsequenceCondSet));
 
   ConditionSets.push_back(isl_set_coalesce(
       isl_set_intersect(ConsequenceCondSet, isl_set_copy(Domain))));
   ConditionSets.push_back(isl_set_coalesce(
       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<BranchInst>(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<ICmpInst>(Cond);
-  ConstantInt *CondConstant = dyn_cast<ConstantInt>(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<SCEVCouldNotCompute>(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);
 }
 
 void ScopStmt::deriveAssumptionsFromGEP(GetElementPtrInst *GEP) {
   int Dimension = 0;
   isl_ctx *Ctx = Parent.getIslCtx();
   isl_local_space *LSpace = isl_local_space_from_space(getDomainSpace());
void Scop::buildDomains(Region *R, LoopInfo &LI, ScopDetection &SD,
                         DominatorTree &DT) {
 
   auto *EntryBB = R->getEntry();
   int LD = getRelativeLoopDepth(LI.getLoopFor(EntryBB));
   auto *S = isl_set_universe(isl_space_set_alloc(getIslCtx(), 0, LD + 1));
   DomainMap[EntryBB] = S;
 
   buildDomainsWithBranchConstraints(R, LI, SD, DT);
+  addLoopBoundsToHeaderDomains(LI, SD, DT);
+  propagateDomainConstraints(R, LI, SD, DT);
 }
 
 void Scop::buildDomainsWithBranchConstraints(Region *R, LoopInfo &LI,
                                              ScopDetection &SD,
                                              DominatorTree &DT) {
   RegionInfo &RI = *R->getRegionInfo();
 
   // To create the domain for each block in R we iterate over all blocks and
for (unsigned u = 0, e = ConditionSets.size(); u < e; u++) {
 
       SuccDomain = isl_set_coalesce(SuccDomain);
       DEBUG(dbgs() << "\tSet SuccBB: " << SuccBB->getName() << " : " << Domain
                    << "\n");
     }
   }
 }
 
+/// @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<BasicBlock *, isl_set *> &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<Region *> 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<Region>();
+      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<Loop *> &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<Loop *, 8> &RegionLoops,
+                     const ScopDetection::BoxedLoopsSetTy &BoxedLoops) {
+
+  SmallVector<Loop *, 8> 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<Loop *, 8> 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<BranchInst>(LatchBB->getTerminator());
+    if (BI->isUnconditional())
+      BackedgeCondition = isl_set_copy(LatchBBDom);
+    else {
+      SmallVector<isl_set *, 2> 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;
 
   if (buildAliasGroups(AA))
     return;
 
   // If a problem occurs while building the alias groups we need to delete

polly/trunk/test/Isl/CodeGen/OpenMP/loop-body-references-outer-values-3.ll

 ; RUN: opt %loadPolly -polly-detect-unprofitable -basicaa -polly-parallel -polly-parallel-force -polly-ast -analyze < %s | FileCheck %s -check-prefix=AST
 ; RUN: opt %loadPolly -polly-detect-unprofitable -basicaa -polly-parallel -polly-parallel-force -polly-codegen -S -verify-dom-info < %s | FileCheck %s -check-prefix=IR
 ; RUN: opt %loadPolly -polly-detect-unprofitable -basicaa -polly-parallel -polly-parallel-force -polly-codegen -S -verify-dom-info < %s | FileCheck %s -check-prefix=IR
 
 ; The interesting part of this test case is the instruction:
 ;   %tmp = bitcast i8* %call to i64**
 ; which is not part of the scop. In the SCEV based code generation not '%tmp',
 ; 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
 
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 define void @foo(i64 %cols, i8* noalias %call) {
 entry:

 
 end:
   ret void
 }
 
 ; 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
 
 define void @bar(i64 %cols, i8* noalias %call) {
 entry:
   %tmp = bitcast i8* %call to i64**
   %arrayidx = getelementptr inbounds i64*, i64** %tmp, i64 0

polly/trunk/test/Isl/CodeGen/OpenMP/reference-preceeding-loop.ll

 ; RUN: opt %loadPolly -polly-detect-unprofitable -polly-parallel -polly-parallel-force -polly-ast -analyze < %s | FileCheck %s -check-prefix=AST
 ; RUN: opt %loadPolly -polly-detect-unprofitable -polly-parallel -polly-parallel-force -polly-codegen -S -verify-dom-info < %s | FileCheck %s -check-prefix=IR
 
 
 ; - 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"
 
 @cum_freq = external global [258 x i64], align 16
 
 define void @update_model(i64 %symbol) {
 entry:
   br label %for.one

polly/trunk/test/Isl/CodeGen/OpenMP/two-parallel-loops-reference-outer-indvar.ll

 ; RUN: opt %loadPolly -polly-detect-unprofitable -polly-parallel -polly-parallel-force -polly-ast -analyze < %s | FileCheck %s -check-prefix=AST
 ; RUN: opt %loadPolly -polly-detect-unprofitable -polly-parallel -polly-parallel-force -polly-codegen -S -verify-dom-info < %s | FileCheck %s -check-prefix=IR
 
 ; 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"
 
 define void @foo(i64 %nj, [512 x double]* %R) {
 entry:
   br label %for.cond1.preheader
 
 for.cond1.preheader:

polly/trunk/test/Isl/CodeGen/phi_scalar_simple_1.ll

 ; CHECK-DAG:     %x.addr.1.s2a = alloca i32
 ; CHECK-DAG:     %x.addr.1.phiops = alloca i32
 ; CHECK-DAG:     %x.addr.0.s2a = alloca i32
 ; CHECK-DAG:     %x.addr.0.phiops = alloca i32
   %tmp = sext i32 %N to i64
   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
 ; 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
   %indvars.iv = phi i64 [ %indvars.iv.next, %for.inc4 ], [ 1, %entry ]
   %x.addr.0 = phi i32 [ %x, %entry ], [ %x.addr.1.lcssa, %for.inc4 ]
   %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
 ; CHECK:         store i32 %x.addr.1.phiops.reload, i32* %x.addr.1.s2a[[R6:[0-9]*]]
 ; CHECK:         store i32 %x.addr.1.phiops.reload, i32* %x.addr.1.lcssa.phiops
   %x.addr.1 = phi i32 [ %x.addr.0, %for.body ], [ %add, %for.inc ]
   %j.0 = phi i32 [ 3, %for.body ], [ %inc, %for.inc ]
   %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
 ; CHECK:         %p_add = add nsw i32 %x.addr.1.s2a.reload[[R3]], %tmp1_p_scalar_
 ; CHECK:         store i32 %p_add, i32* %x.addr.1.phiops
   %arrayidx = getelementptr inbounds i32, i32* %A, i64 %indvars.iv
   %tmp1 = load i32, i32* %arrayidx, align 4
   %add = add nsw i32 %x.addr.1, %tmp1
   %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
 }

polly/trunk/test/Isl/CodeGen/two-scops-in-row.ll

 
 ; RUN: opt %loadPolly -polly-detect-unprofitable -polly-no-early-exit -polly-ast -analyze -polly-ignore-aliasing < %s | FileCheck %s -check-prefix=SCALAR
 ; RUN: opt %loadPolly -polly-detect-unprofitable -polly-no-early-exit -polly-codegen -polly-ignore-aliasing -disable-output < %s
 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(c0);
-; SCALAR:       if (Scalar0.val >= 100)
-; SCALAR:         Stmt_for_1(0);
+; SCALAR:       Stmt_for_1(0);
+; SCALAR:       for (int c0 = 1; c0 <= -Scalar0.val + 99; c0 += 1)
+; SCALAR:         Stmt_for_1(c0);
 ; SCALAR:     }
 
 ; SCALAR: if (1)
 ; SCALAR:     Stmt_for_0(0);
 
 
 define void @foo(i32* %A) {
 entry:

polly/trunk/test/ScopInfo/NonAffine/div_backedge.ll

+; 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
+}

polly/trunk/test/ScopInfo/NonAffine/div_domain.ll

+; 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
+}

polly/trunk/test/ScopInfo/NonAffine/modulo_backedge.ll

+; 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
+}

polly/trunk/test/ScopInfo/NonAffine/modulo_domain.ll

+; 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
+}

polly/trunk/test/ScopInfo/NonAffine/non-affine-loop-condition-dependent-access_2.ll

 ; innermost loop as a SCoP of depth 1, we can overapproximate the
 ; innermost loop in the whole loop nest and model A[k] as a non-affine
 ; access.
 ;
 ; INNERMOST:    Function: f
 ; INNERMOST:    Region: %bb15---%bb13
 ; INNERMOST:    Max Loop Depth:  1
 ; INNERMOST:    p0: {0,+,{0,+,1}<nuw><nsw><%bb11>}<nuw><nsw><%bb13>
-; INNERMOST:    p1: {0,+,{0,+,-1}<nw><%bb11>}<nw><%bb13>
-; INNERMOST:    p2: {0,+,4}<nuw><nsw><%bb11>
-; INNERMOST:    p3: {0,+,4}<nuw><nsw><%bb13>
-; INNERMOST:    p4: {0,+,{0,+,4}<nuw><nsw><%bb11>}<%bb13>
+; INNERMOST:    p1: {0,+,4}<nuw><nsw><%bb11>
+; INNERMOST:    p2: {0,+,4}<nuw><nsw><%bb13>
+; INNERMOST:    p3: {0,+,{0,+,4}<nuw><nsw><%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
 ; ALL:    Region: %bb11---%bb29
 ; ALL:    Max Loop Depth:  2
 ; ALL:    Context:
 ; ALL:    {  :  }
 ; ALL:    Assumed Context:

polly/trunk/test/ScopInfo/NonAffine/non-affine-loop-condition-dependent-access_3.ll

 ; INNERMOST:            ReadAccess := [Reduction Type: NONE] [Scalar: 0]
 ; INNERMOST:                [p_0, p_1, p_2] -> { Stmt_bb16[i0] -> MemRef_A[o0] : 4o0 = p_2 };
 ; INNERMOST:            ReadAccess := [Reduction Type: +] [Scalar: 0]
 ; INNERMOST:                [p_0, p_1, p_2] -> { Stmt_bb16[i0] -> MemRef_A[i0] };
 ; INNERMOST:            MustWriteAccess :=  [Reduction Type: +] [Scalar: 0]
 ; 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]
 ; INNERMOST:                [p_0, p_1, p_2] -> { Stmt_bb26[] -> MemRef_indvars_iv_next6[] };
 ; INNERMOST:            MustWriteAccess :=  [Reduction Type: NONE] [Scalar: 1]
 ; INNERMOST:                [p_0, p_1, p_2] -> { Stmt_bb26[] -> MemRef_indvars_iv_next4[] };
 ; INNERMOST:    }
 ;

polly/trunk/test/ScopInfo/isl_trip_count_01.ll

 ; 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;
 ;     while (M <= N) {
 ;       A[i++] = 1;
 ;       M += 4;
 ;     }
 ;   }

polly/trunk/test/ScopInfo/loop_affine_bound_0.ll

 ; CHECK:              [N, M] -> { Stmt_bb1[i0, i1] :
 ; CHECK-DAG:             i0 >= 0
 ; CHECK-DAG:          and
 ; CHECK-DAG:             i0 <= 2 + 4N + 7M
 ; CHECK-DAG:          and
 ; 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] };
 ; CHECK:          MustWriteAccess := [Reduction Type: NONE]
 ; CHECK:              [N, M] -> { Stmt_bb1[i0, i1] -> MemRef_a[i0 + 128i1] };
 ; CHECK:  }

polly/trunk/test/ScopInfo/loop_affine_bound_1.ll

 
 define void @f([128 x i64]* nocapture %a, i64 %N, i64 %M) nounwind {
 entry:
   %0 = shl i64 %N, 2                              ; <i64> [#uses=2]
   %1 = mul i64 %M, 7                              ; <i64> [#uses=2]
   %2 = or i64 %0, 3                               ; <i64> [#uses=1]
   %3 = add nsw i64 %2, %1                         ; <i64> [#uses=1]
   %4 = icmp sgt i64 %3, 0                         ; <i1> [#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 ] ; <i64> [#uses=2]
   %scevgep = getelementptr [128 x i64], [128 x i64]* %a, i64 %indvar, i64 %tmp10 ; <i64*> [#uses=1]
   store i64 0, i64* %scevgep, align 8
   %indvar.next = add i64 %indvar, 1               ; <i64> [#uses=2]
   %exitcond = icmp sge i64 %indvar.next, %tmp9     ; <i1> [#uses=1]
   br i1 %exitcond, label %bb3, label %bb1

 }
 
 ; CHECK: p0: %N
 ; CHECK: p1: %M
 ; CHECK: Statements {
 ; 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: }

polly/trunk/test/ScopInfo/loop_affine_bound_2.ll

 ; CHECK:             [N, M] -> { Stmt_bb1[i0, i1] :
 ; CHECK-DAG:                  i0 >= 0
 ; CHECK-DAG:                and
 ; CHECK-DAG:                  i0 <= 2 + 4N + 7M
 ; CHECK-DAG:                and
 ; 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]
 ; CHECK:         MustWriteAccess := [Reduction Type: NONE]
 ; CHECK:             [N, M] -> { Stmt_bb1[i0, i1] -> MemRef_a[-1152 + 768M + 897i0 + 128i1] };
 ; CHECK: }

polly/trunk/test/ScopInfo/non_affine_region_1.ll

 ; CHECK:       Stmt_bb10__TO__bb18
 ; CHECK-NEXT:        Domain :=
 ; CHECK-NEXT:            [b] -> { Stmt_bb10__TO__bb18[i0] :
 ; CHECK-DAG:               i0 >= 0
 ; CHECK-DAG:             and
 ; 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]
 ;
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 define void @f(i32* %A, i32 %b) {
 bb:
   br label %bb1

polly/trunk/test/ScopInfo/pointer-type-expressions.ll

 }
 
 ; CHECK: Assumed Context:
 ; CHECK:   {  :  }
 
 ; 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 };
 ; CHECK:        MustWriteAccess := [Reduction Type: NONE]
 ; CHECK:            [P, N] -> { Stmt_store[i0] -> MemRef_a[i0] };

polly/trunk/test/ScopInfo/redundant_parameter_constraint.ll

+; 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
+}

polly/trunk/test/ScopInfo/two-loops-right-after-each-other.ll

 ; RUN: opt %loadPolly -polly-detect-unprofitable -polly-scops -analyze \
 ; RUN: < %s | FileCheck %s
 
 ; 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:
   br label %branch
 
 branch:
   %cond = icmp sle i64 %N, 100
   br i1 %cond, label %loop.1, label %merge

polly/trunk/test/ScopInfo/unsigned-condition.ll

 }
 
 ; CHECK: Assumed Context:
 ; CHECK:   {  :  }
 
 ; 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] };
 ; CHECK:        MustWriteAccess := [Reduction Type: NONE]
 ; CHECK:            [P, N] -> { Stmt_store[i0] -> MemRef_a[i0] };