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

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

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

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

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

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

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

test/ScopInfo/loop_affine_bound_1.ll

19

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

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

test/ScopInfo/loop_affine_bound_2.ll

77

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

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

912

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

lib/Analysis/ScopInfo.cpp

817

bset, maybe uppercase?

828

Please comment the contents of the return value.

831

Uppercase iterators?

841

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

Unnecessary line break

857

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

1705

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

Why an empty line?

1753

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

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

799

artificial

1434

Function extraction into separate commit?

1536

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

1575

Empty comment line/missing annotations

1675

hence -> e.g.

1760

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

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

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 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.
   /// @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);

lib/Analysis/ScopInfo.cpp

 //===--------- ScopInfo.cpp  - Create Scops from LLVM IR ------------------===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 //

 #include "polly/TempScopInfo.h"
 #include "llvm/ADT/MapVector.h"
 #include "llvm/ADT/SetVector.h"
 #include "llvm/ADT/Statistic.h"
 #include "llvm/ADT/STLExtras.h"
 #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"
 #include "llvm/Support/Debug.h"
 #include "isl/aff.h"
 #include "isl/constraint.h"
 #include "isl/local_space.h"
 #include "isl/map.h"
if (CCond->isZero())
       ConsequenceCondSet = isl_set_empty(isl_set_get_space(Domain));
     else
       ConsequenceCondSet = isl_set_universe(isl_set_get_space(Domain));
   } else {
     auto *ICond = dyn_cast<ICmpInst>(Condition);
     assert(ICond &&
            "Condition of exiting branch was neither constant nor ICmp!");
 
     ScalarEvolution &SE = *S.getSE();
     isl_pw_aff *LHS, *RHS;
     LHS = S.getPwAff(SE.getSCEVAtScope(ICond->getOperand(0), L), Domain);
     RHS = S.getPwAff(SE.getSCEVAtScope(ICond->getOperand(1), L), Domain);
     ConsequenceCondSet = buildConditionSet(ICond->getPredicate(), LHS, RHS);
   }
 
   assert(ConsequenceCondSet);
   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");
-  }
+/// @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<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;
 
-  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());
+/// @brief Return @p S without basic sets that are unbounded in dimension @p
+/// Dim.
+static __isl_give std::pair<isl_set *, isl_set *>
+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_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);
 
-  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.
-      addLoopTripCountToDomain(L);
-      isl_pw_aff_free(IV);
-    }
-  }
-
-  isl_local_space_free(LocalSpace);
+  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);
 }
 
 void ScopStmt::deriveAssumptionsFromGEP(GetElementPtrInst *GEP) {
   int Dimension = 0;
   isl_ctx *Ctx = Parent.getIslCtx();
   isl_local_space *LSpace = isl_local_space_from_space(getDomainSpace());
static inline BasicBlock *getRegionNodeSuccessor(RegionNode *RN, BranchInst *BI,
                                                  unsigned idx) {
   if (RN->isSubRegion()) {
     assert(idx == 0);
     return RN->getNodeAs<Region>()->getExit();
   }
   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<BasicBlock>());
 
   Region *NonAffineSubRegion = RN->getNodeAs<Region>();
-  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();
   return isl_set_copy(DomainMap[BB]);
 }
 
 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);
+  addLoopConstraintsToDomains(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
if (RN->isSubRegion()) {
       if (!SD.isNonAffineSubRegion(SubRegion, &getRegion())) {
         buildDomainsWithBranchConstraints(SubRegion, LI, SD, DT);
         continue;
       }
     }
 
     BasicBlock *BB = getRegionNodeBasicBlock(RN);
     isl_set *Domain = DomainMap[BB];
     DEBUG(dbgs() << "\tVisit: " << BB->getName() << " : " << Domain << "\n");
     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.");
 
     Loop *BBLoop = getRegionNodeLoop(RN, LI);
     int BBLoopDepth = getRelativeLoopDepth(BBLoop);
 
for (unsigned u = 0, e = ConditionSets.size(); u < e; u++) {
 
       SuccDomain = isl_set_coalesce(SuccDomain);
       DEBUG(dbgs() << "\tSet SuccBB: " << SuccBB->getName() << " : " << Domain
                    << "\n");
     }
   }
 }
 
+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) {
+
+  const ScopDetection::BoxedLoopsSetTy *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];
+    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<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 all 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());
+  }
+}
+
+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<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);
+      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;
 
   if (buildAliasGroups(AA))
     return;
 
   // If a problem occurs while building the alias groups we need to delete

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

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

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:

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

 ; and also a non-affine access (A[k]). While we can always model the
 ; 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---%bb26
 ; INNERMOST:    Max Loop Depth:  1
+;OLDINNERMOST:    p0: {0,+,{0,+,1}<nuw><nsw><%bb11>}<nuw><nsw><%bb13>
+;OLDINNERMOST:    p1: {0,+,{0,+,-1}<nw><%bb11>}<nw><%bb13>
+;OLDINNERMOST:    p2: {0,+,4}<nuw><nsw><%bb11>
+;OLDINNERMOST:    p3: {0,+,4}<nuw><nsw><%bb13>
+;OLDINNERMOST:    p4: {0,+,{0,+,4}<nuw><nsw><%bb11>}<%bb13>
 ; 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 }
+;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
 ; ALL:    Region: %bb11---%bb29
 ; ALL:    Max Loop Depth:  2
 ; ALL:    Context:
 ; ALL:    {  :  }
 ; ALL:    Assumed Context:

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

test/ScopInfo/isl_trip_count_02.ll

 ; 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++)
 ;       A[i] = i;
 ;   }
 ;
 target datalayout = "e-m:e-p:32:32-i64:64-v128:64:128-a:0:32-n32-S64"
 

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:  }

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: }

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: }

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

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

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