[Polly] Track assumptions and restrictions separately

[jdoerfert]

Summary

In order to speed up compile time and to avoid random timeouts we now
separately track assumptions and restrictions. In this context assumptions
describe parameter valuations we need and restrictions describe parameter
valuations we do not allow. During AST generation we create a runtime check
for both, whereas the one for the restrictions is negated before a
conjunction is build.

Except the In-Bounds assumptions we currently only track restrictions.

[grosser]

Hi Johannes,

thanks for this patch. It looks very nice. I just have two semantical issues, I need feedback on and two smaller stylistic comments.

include/polly/ScopInfo.h

1313

Comment is incomplete

lib/Analysis/ScopInfo.cpp

591

Is there a specific reason we do not drop the complement and use the other assumption context?

594

Plain boolean parameters make it difficult to understand the code without knowing the name of parameter 4. Maybe just introduce an enum that clearly describes the two cases? Maybe ASSUME_VALID, ASSUME_INVALID?

1312

Nice.

1748

Nice.

1913

Did you check that the very same simplication logic also works for invalid context.

Assuming you have code

   for (i = 0; i < N; i++)
S:     ...

here we have N > 0 when we project the domain on the parameter set.

Now, assuming the code is invalid for N > 1000. If we gist_simplify the invalid set, we loose this constraint, no?

2406

Nice to drop the complement.

2507

Nice to drop the complement.

3093

This should not be NULL, no?

3103

This should not be NULL, no?

lib/CodeGen/IslAst.cpp

337

Can we only add the condition of the invalid context if it is non-trivial? This will make the output nicer and also should reduce the number of test cases changed?

test/Isl/CodeGen/exprModDiv.ll

42

This change seems to be independently beneficial. Maybe it could be committed ahead of time?

test/ScopInfo/assume_gep_bounds_2.ll

19

Unnecessary whitespace changes?

test/ScopInfo/long-sequence-of-error-blocks-2.ll

10

Nice!

[jdoerfert]

Update according to the comments

lib/Analysis/ScopInfo.cpp

591

This way all INBOUNDS assumption are "positve" thus assumptions. We could make it an restriction here easily but the current implementation of INBOUNDS further down will always need a complement and the current choice is probably more effiecent.

Tl;dr
No "real" reason. It just looks weird to have INBOUNDS assumptions as well as restrictions...

1913

Mh,... I removed the secon call for now. I have to think about this.

[grosser]

LGTM. Thank you Johannes!

include/polly/ScopInfo.h

private:
   ///
   /// When constructing a scop sometimes the exact representation of a statement
   /// or condition would be very complex, but there is a common case which is a
   /// lot simpler, but which is only valid under certain assumptions. The
   /// assumed context records the assumptions taken during the construction of
   /// this scop and that need to be code generated as a run-time test.
   isl_set *AssumedContext;
 
-  /// @brief The boundary assumptions under which this scop was built.
+  /// @brief The restrictions under which this scop was built.
   ///
-  /// The boundary context is similar to the assumed context as it contains
+  /// The invalid context is similar to the assumed context as it contains
   /// constraints over the parameters we assume to be true. However, the
-  /// boundary context is less useful for dependence analysis and
-  /// simplification purposes as it contains only constraints that affect the
-  /// boundaries of the parameter ranges. As these constraints can become quite
-  /// complex, the boundary context and the assumed context are separated as a
-  /// meassure to save compile time.
-  isl_set *BoundaryContext;
+  /// invalid context ... TODO
+  isl_set *InvalidContext;
 
   /// @brief The schedule of the SCoP
   ///
   /// The schedule of the SCoP describes the execution order of the statements
   /// in the scop by assigning each statement instance a possibly
   /// multi-dimensional execution time. The schedule is stored as a tree of
   /// schedule nodes.
   ///
private:
   void hoistInvariantLoads(ScopDetection &SD);
 
   /// @brief Add invariant loads listed in @p InvMAs with the domain of @p Stmt.
   void addInvariantLoads(ScopStmt &Stmt, MemoryAccessList &InvMAs);
 
   /// @brief Build the Context of the Scop.
   void buildContext();
 
-  /// @brief Build the BoundaryContext based on the wrapping of expressions.
-  void buildBoundaryContext();
+  /// @brief Add the restrictions based on the wrapping of expressions.
+  void addWrappingContext();
 
   /// @brief Add user provided parameter constraints to context (source code).
   void addUserAssumptions(AssumptionCache &AC, DominatorTree &DT, LoopInfo &LI);
 
   /// @brief Add user provided parameter constraints to context (command line).
   void addUserContext();
 
   /// @brief Add the bounds of the parameters to the context.
   void addParameterBounds();
 
-  /// @brief Simplify the assumed and boundary context.
+  /// @brief Simplify the assumed and invalid context.
   void simplifyContexts();
 
   /// @brief Get the representing SCEV for @p S if applicable, otherwise @p S.
   ///
   /// Invariant loads of the same location are put in an equivalence class and
   /// only one of them is chosen as a representing element that will be
   /// modeled as a parameter. The others have to be normalized, i.e.,
   /// replaced by the representing element of their equivalence class, in order
public:
   __isl_give isl_set *getContext() const;
   __isl_give isl_space *getParamSpace() const;
 
   /// @brief Get the assumed context for this Scop.
   ///
   /// @return The assumed context of this Scop.
   __isl_give isl_set *getAssumedContext() const;
 
-  /// @brief Get the runtime check context for this Scop.
-  ///
-  /// The runtime check context contains all constraints that have to
-  /// hold at runtime for the optimized version to be executed.
-  ///
-  /// @return The runtime check context of this Scop.
-  __isl_give isl_set *getRuntimeCheckContext() const;
-
   /// @brief Return true if the optimized SCoP can be executed.
   ///
   /// In addition to the runtime check context this will also utilize the domain
   /// constraints to decide it the optimized version can actually be executed.
   ///
   /// @returns True if the optimized SCoP can be executed.
   bool hasFeasibleRuntimeContext() const;
 
   /// @brief Track and report an assumption.
   ///
-  /// Use 'clang -Rpass-analysis=polly-scops' or 'opt -pass-remarks=polly-scops'
-  /// to output the assumptions.
+  /// Use 'clang -Rpass-analysis=polly-scops' or 'opt
+  /// -pass-remarks-analysis=polly-scops' to output the assumptions.
   ///
-  /// @param Kind The assumption kind describing the underlying cause.
-  /// @param Set  The relations between parameters that are assumed to hold.
-  /// @param Loc  The location in the source that caused this assumption.
-  void trackAssumption(AssumptionKind Kind, __isl_keep isl_set *Set,
-                       DebugLoc Loc);
+  /// @param Kind  The assumption kind describing the underlying cause.
+  /// @param Set   The relations between parameters that are assumed to hold.
+  /// @param Loc   The location in the source that caused this assumption.
+  /// @param IsPos Flag to indicate if the assumptions in @p Set are positive
+  ///              (needed) or negative (invalid).
+  ///
+  /// @returns True if the assumption is not trivial.
+  bool trackAssumption(AssumptionKind Kind, __isl_keep isl_set *Set,
+                       DebugLoc Loc, bool IsPos);
 
   /// @brief Add assumptions to assumed context.
   ///
   /// The assumptions added will be assumed to hold during the execution of the
   /// scop. However, as they are generally not statically provable, at code
   /// generation time run-time checks will be generated that ensure the
   /// assumptions hold.
   ///
   /// WARNING: We currently exploit in simplifyAssumedContext the knowledge
   ///          that assumptions do not change the set of statement instances
   ///          executed.
   ///
   /// @param Kind The assumption kind describing the underlying cause.
   /// @param Set  The relations between parameters that are assumed to hold.
   /// @param Loc  The location in the source that caused this assumption.
-  void addAssumption(AssumptionKind Kind, __isl_take isl_set *Set,
-                     DebugLoc Loc);
+  /// @param IsPos Flag to indicate if the assumptions in @p Set are positive
+  ///              (needed) or negative (invalid).
+  void addAssumption(AssumptionKind Kind, __isl_take isl_set *Set, DebugLoc Loc,
+                     bool IsPos);
 
   /// @brief Mark the scop as invalid.
   ///
   /// This method adds an assumption to the scop that is always invalid. As a
   /// result, the scop will not be optimized later on. This function is commonly
   /// called when a condition makes it impossible (or too compile time
   /// expensive) to process this scop any further.
   ///
   /// @param Kind The assumption kind describing the underlying cause.
   /// @param Loc  The location in the source that triggered .
   void invalidate(AssumptionKind Kind, DebugLoc Loc);
 
-  /// @brief Get the boundary context for this Scop.
+  /// @brief Get the invalid context for this Scop.
   ///
-  /// @return The boundary context of this Scop.
-  __isl_give isl_set *getBoundaryContext() const;
+  /// @return The invalid context of this Scop.
+  __isl_give isl_set *getInvalidContext() const;
 
   /// @brief Build the alias checks for this SCoP.
   void buildAliasChecks(AliasAnalysis &AA);
 
   /// @brief Build all alias groups for this SCoP.
   ///
   /// @returns True if __no__ error occurred, false otherwise.
   bool buildAliasGroups(AliasAnalysis &AA);
 
   /// @brief Return all alias groups for this SCoP.
   const MinMaxVectorPairVectorTy &getAliasGroups() const {
     return MinMaxAliasGroups;
   }
 
   /// @brief Get an isl string representing the context.
   std::string getContextStr() const;
 
   /// @brief Get an isl string representing the assumed context.
   std::string getAssumedContextStr() const;
 
-  /// @brief Get an isl string representing the boundary context.
-  std::string getBoundaryContextStr() const;
+  /// @brief Get an isl string representing the invalid context.
+  std::string getInvalidContextStr() const;
 
   /// @brief Return the stmt for the given @p BB or nullptr if none.
   ScopStmt *getStmtForBasicBlock(BasicBlock *BB) const;
 
   /// @brief Return the ScopStmt that represents @p RN; can return nullptr if
   /// the RegionNode is not within the SCoP or has been removed due to
   /// simplifications.
   ScopStmt *getStmtForRegionNode(RegionNode *RN) const;

lib/Analysis/ScopInfo.cpp

void MemoryAccess::assumeNoOutOfBound() {
   Outside = isl_set_intersect(Outside, Statement->getDomain());
   Outside = isl_set_params(Outside);
 
   // Remove divs to avoid the construction of overly complicated assumptions.
   // Doing so increases the set of parameter combinations that are assumed to
   // not appear. This is always save, but may make the resulting run-time check
   // bail out more often than strictly necessary.
   Outside = isl_set_remove_divs(Outside);
   Outside = isl_set_complement(Outside);
-  Statement->getParent()->addAssumption(
-      INBOUNDS, Outside,
-      getAccessInstruction() ? getAccessInstruction()->getDebugLoc() : nullptr);
+  auto &Loc = getAccessInstruction() ? getAccessInstruction()->getDebugLoc()
+                                     : DebugLoc();
+  Statement->getParent()->addAssumption(INBOUNDS, Outside, Loc, true);
   isl_space_free(Space);
 }
 
 void MemoryAccess::computeBoundsOnAccessRelation(unsigned ElementSize) {
   ScalarEvolution *SE = Statement->getParent()->getSE();
 
   Value *Ptr = MemAccInst(getAccessInstruction()).getPointerOperand();
   if (!Ptr || !SE->isSCEVable(Ptr->getType()))
void ScopStmt::deriveAssumptionsFromGEP(GetElementPtrInst *GEP,
   if (auto *PtrTy = dyn_cast<PointerType>(Ty)) {
     Ty = PtrTy->getElementType();
   }
 
   int IndexOffset = Subscripts.size() - Sizes.size();
 
   assert(IndexOffset <= 1 && "Unexpected large index offset");
 
+  auto *NotExecuted = isl_set_complement(isl_set_params(getDomain()));
   for (size_t i = 0; i < Sizes.size(); i++) {
     auto Expr = Subscripts[i + IndexOffset];
     auto Size = Sizes[i];
 
     InvariantLoadsSetTy AccessILS;
     if (!isAffineExpr(&Parent.getRegion(), Expr, SE, nullptr, &AccessILS))
       continue;
 
for (size_t i = 0; i < Sizes.size(); i++) {
 
     isl_pw_aff *DimSize = isl_pw_aff_from_aff(isl_aff_val_on_domain(
         isl_local_space_copy(LSpace), isl_val_int_from_si(Ctx, Size)));
 
     isl_set *OutOfBound = isl_pw_aff_ge_set(AccessOffset, DimSize);
     OutOfBound = isl_set_intersect(getDomain(), OutOfBound);
     OutOfBound = isl_set_params(OutOfBound);
     isl_set *InBound = isl_set_complement(OutOfBound);
-    isl_set *Executed = isl_set_params(getDomain());
 
     // A => B == !A or B
     isl_set *InBoundIfExecuted =
-        isl_set_union(isl_set_complement(Executed), InBound);
+        isl_set_union(isl_set_copy(NotExecuted), InBound);
 
     InBoundIfExecuted = isl_set_coalesce(InBoundIfExecuted);
-    Parent.addAssumption(INBOUNDS, InBoundIfExecuted, GEP->getDebugLoc());
+    Parent.addAssumption(INBOUNDS, InBoundIfExecuted, GEP->getDebugLoc(), true);
   }
 
   isl_local_space_free(LSpace);
+  isl_set_free(NotExecuted);
 }
 
 void ScopStmt::deriveAssumptions(BasicBlock *Block, ScopDetection &SD) {
   for (Instruction &Inst : *Block)
     if (auto *GEP = dyn_cast<GetElementPtrInst>(&Inst))
       deriveAssumptionsFromGEP(GEP, SD);
 }
 
return isl_id_alloc(getIslCtx(), ParameterName.c_str(),
                       const_cast<void *>((const void *)Parameter));
 }
 
 isl_set *Scop::addNonEmptyDomainConstraints(isl_set *C) const {
   isl_set *DomainContext = isl_union_set_params(getDomains());
   return isl_set_intersect_params(C, DomainContext);
 }
 
-void Scop::buildBoundaryContext() {
-  if (IgnoreIntegerWrapping) {
-    BoundaryContext = isl_set_universe(getParamSpace());
+void Scop::addWrappingContext() {
+  if (IgnoreIntegerWrapping)
     return;
-  }
-
-  BoundaryContext = Affinator.getWrappingContext();
 
-  // The isl_set_complement operation used to create the boundary context
-  // can possibly become very expensive. We bound the compile time of
-  // this operation by setting a compute out.
-  //
-  // TODO: We can probably get around using isl_set_complement and directly
-  // AST generate BoundaryContext.
-  long MaxOpsOld = isl_ctx_get_max_operations(getIslCtx());
-  isl_ctx_reset_operations(getIslCtx());
-  isl_ctx_set_max_operations(getIslCtx(), 300000);
-  isl_options_set_on_error(getIslCtx(), ISL_ON_ERROR_CONTINUE);
-
-  BoundaryContext = isl_set_complement(BoundaryContext);
-
-  if (isl_ctx_last_error(getIslCtx()) == isl_error_quota) {
-    isl_set_free(BoundaryContext);
-    BoundaryContext = isl_set_empty(getParamSpace());
-  }
-
-  isl_options_set_on_error(getIslCtx(), ISL_ON_ERROR_ABORT);
-  isl_ctx_reset_operations(getIslCtx());
-  isl_ctx_set_max_operations(getIslCtx(), MaxOpsOld);
-  BoundaryContext = isl_set_gist_params(BoundaryContext, getContext());
-  trackAssumption(WRAPPING, BoundaryContext, DebugLoc());
+  auto *WrappingContext = Affinator.getWrappingContext();
+  addAssumption(WRAPPING, WrappingContext, DebugLoc(), false);
 }
 
 void Scop::addUserAssumptions(AssumptionCache &AC, DominatorTree &DT,
                               LoopInfo &LI) {
   auto *R = &getRegion();
   auto &F = *R->getEntry()->getParent();
   for (auto &Assumption : AC.assumptions()) {
     auto *CI = dyn_cast_or_null<CallInst>(Assumption);
for (LoadInst *LInst : RIL) {
     InvariantEquivClasses.emplace_back(PointerSCEV, MemoryAccessList(), nullptr,
                                        Ty);
   }
 }
 
 void Scop::buildContext() {
   isl_space *Space = isl_space_params_alloc(IslCtx, 0);
   Context = isl_set_universe(isl_space_copy(Space));
+  InvalidContext = isl_set_empty(isl_space_copy(Space));
   AssumedContext = isl_set_universe(Space);
 }
 
 void Scop::addParameterBounds() {
   for (const auto &ParamID : ParameterIds) {
     int dim = ParamID.second;
 
     ConstantRange SRange = SE->getSignedRange(ParamID.first);
void Scop::simplifyContexts() {
   //     for (long i = 0; i < 100; i++)
   //       for (long j = 0; j < m; j++)
   //         A[i+p][j] = 1.0;
   //
   //   we assume that the condition m <= 0 or (m >= 1 and p >= 0) holds as
   //   otherwise we would access out of bound data. Now, knowing that code is
   //   only executed for the case m >= 0, it is sufficient to assume p >= 0.
   AssumedContext = simplifyAssumptionContext(AssumedContext, *this);
-  BoundaryContext = simplifyAssumptionContext(BoundaryContext, *this);
+  InvalidContext = simplifyAssumptionContext(InvalidContext, *this);
 }
 
 /// @brief Add the minimal/maximal access in @p Set to @p User.
 static isl_stat buildMinMaxAccess(__isl_take isl_set *Set, void *User) {
   Scop::MinMaxVectorTy *MinMaxAccesses = (Scop::MinMaxVectorTy *)User;
   isl_pw_multi_aff *MinPMA, *MaxPMA;
   isl_pw_aff *LastDimAff;
   isl_aff *OneAff;
for (auto *RN : RTraversal) {
 
     if (BBLoop && BBLoop->getHeader() == BB && getRegion().contains(BBLoop))
       addLoopBoundsToHeaderDomain(BBLoop, LI);
 
     // Add assumptions for error blocks.
     if (containsErrorBlock(RN, getRegion(), LI, DT)) {
       IsOptimized = true;
       isl_set *DomPar = isl_set_params(isl_set_copy(Domain));
-      addAssumption(ERRORBLOCK, isl_set_complement(DomPar),
-                    BB->getTerminator()->getDebugLoc());
+      addAssumption(ERRORBLOCK, DomPar, BB->getTerminator()->getDebugLoc(),
+                    false);
     }
   }
 }
 
 /// @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.
void Scop::addLoopBoundsToHeaderDomain(Loop *L, LoopInfo &LI) {
   // the bounded assumptions to the context as they are already implied by the
   // <nsw> tag.
   if (Affinator.hasNSWAddRecForLoop(L)) {
     isl_set_free(Parts.first);
     return;
   }
 
   isl_set *UnboundedCtx = isl_set_params(Parts.first);
-  isl_set *BoundedCtx = isl_set_complement(UnboundedCtx);
-  addAssumption(INFINITELOOP, BoundedCtx,
-                HeaderBB->getTerminator()->getDebugLoc());
+  addAssumption(INFINITELOOP, UnboundedCtx,
+                HeaderBB->getTerminator()->getDebugLoc(), false);
 }
 
 void Scop::buildAliasChecks(AliasAnalysis &AA) {
   if (!PollyUseRuntimeAliasChecks)
     return;
 
   if (buildAliasGroups(AA))
     return;
static unsigned getMaxLoopDepthInRegion(const Region &R, LoopInfo &LI,
   return MaxLD - MinLD + 1;
 }
 
 Scop::Scop(Region &R, ScalarEvolution &ScalarEvolution, isl_ctx *Context,
            unsigned MaxLoopDepth)
     : SE(&ScalarEvolution), R(R), IsOptimized(false),
       HasSingleExitEdge(R.getExitingBlock()), HasErrorBlock(false),
       MaxLoopDepth(MaxLoopDepth), IslCtx(Context), Context(nullptr),
-      Affinator(this), AssumedContext(nullptr), BoundaryContext(nullptr),
+      Affinator(this), AssumedContext(nullptr), InvalidContext(nullptr),
       Schedule(nullptr) {
   buildContext();
 }
 
 void Scop::init(AliasAnalysis &AA, AssumptionCache &AC, ScopDetection &SD,
                 DominatorTree &DT, LoopInfo &LI) {
   addUserAssumptions(AC, DT, LI);
   buildInvariantEquivalenceClasses(SD);
 
   buildDomains(&R, SD, DT, LI);
 
   // Remove empty and ignored statements.
   // Exit early in case there are no executable statements left in this scop.
   simplifySCoP(true, DT, LI);
   if (Stmts.empty())
     return;
 
   // The ScopStmts now have enough information to initialize themselves.
   for (ScopStmt &Stmt : Stmts)
     Stmt.init(SD);
 
   buildSchedule(SD, LI);
 
-  if (isl_set_is_empty(AssumedContext))
+  if (!hasFeasibleRuntimeContext())
     return;
 
   updateAccessDimensionality();
   realignParams();
   addParameterBounds();
   addUserContext();
-  buildBoundaryContext();
+  addWrappingContext();
   simplifyContexts();
   buildAliasChecks(AA);
 
   hoistInvariantLoads(SD);
   simplifySCoP(false, DT, LI);
 }
 
 Scop::~Scop() {
   isl_set_free(Context);
   isl_set_free(AssumedContext);
-  isl_set_free(BoundaryContext);
+  isl_set_free(InvalidContext);
   isl_schedule_free(Schedule);
 
   for (auto It : DomainMap)
     isl_set_free(It.second);
 
   // Free the alias groups
   for (MinMaxVectorPairTy &MinMaxAccessPair : MinMaxAliasGroups) {
     for (MinMaxAccessTy &MMA : MinMaxAccessPair.first) {
const ScopArrayInfo *Scop::getScopArrayInfo(Value *BasePtr,
                                             ScopArrayInfo::MemoryKind Kind) {
   auto *SAI = ScopArrayInfoMap[std::make_pair(BasePtr, Kind)].get();
   assert(SAI && "No ScopArrayInfo available for this base pointer");
   return SAI;
 }
 
 std::string Scop::getContextStr() const { return stringFromIslObj(Context); }
 std::string Scop::getAssumedContextStr() const {
+  assert(AssumedContext && "Assumed context not yet built");
   return stringFromIslObj(AssumedContext);
 }
-std::string Scop::getBoundaryContextStr() const {
-  return stringFromIslObj(BoundaryContext);
+std::string Scop::getInvalidContextStr() const {
+  return stringFromIslObj(InvalidContext);
 }
 
 std::string Scop::getNameStr() const {
   std::string ExitName, EntryName;
   raw_string_ostream ExitStr(ExitName);
   raw_string_ostream EntryStr(EntryName);
 
   R.getEntry()->printAsOperand(EntryStr, false);

 }
 
 __isl_give isl_set *Scop::getContext() const { return isl_set_copy(Context); }
 __isl_give isl_space *Scop::getParamSpace() const {
   return isl_set_get_space(Context);
 }
 
 __isl_give isl_set *Scop::getAssumedContext() const {
+  assert(AssumedContext && "Assumed context not yet built");
   return isl_set_copy(AssumedContext);
 }
 
-__isl_give isl_set *Scop::getRuntimeCheckContext() const {
-  isl_set *RuntimeCheckContext = getAssumedContext();
-  RuntimeCheckContext =
-      isl_set_intersect(RuntimeCheckContext, getBoundaryContext());
-  RuntimeCheckContext = simplifyAssumptionContext(RuntimeCheckContext, *this);
-  return RuntimeCheckContext;
-}
-
 bool Scop::hasFeasibleRuntimeContext() const {
-  isl_set *RuntimeCheckContext = getRuntimeCheckContext();
-  RuntimeCheckContext = addNonEmptyDomainConstraints(RuntimeCheckContext);
-  bool IsFeasible = !isl_set_is_empty(RuntimeCheckContext);
-  isl_set_free(RuntimeCheckContext);
+  isl_set *PositiveContext = getAssumedContext();
+  if (!PositiveContext)
+    return false;
+
+  PositiveContext = addNonEmptyDomainConstraints(PositiveContext);
+  bool IsFeasible = !isl_set_is_empty(PositiveContext);
+  isl_set_free(PositiveContext);
+  if (!IsFeasible)
+    return false;
+
+  isl_set *NegativeContext = getInvalidContext();
+  if (!NegativeContext)
+    return false;
+
+  isl_set *DomainContext = isl_union_set_params(getDomains());
+  IsFeasible = !isl_set_is_subset(DomainContext, NegativeContext);
+  isl_set_free(NegativeContext);
+  isl_set_free(DomainContext);
+
   return IsFeasible;
 }
 
 static std::string toString(AssumptionKind Kind) {
   switch (Kind) {
   case ALIASING:
     return "No-aliasing";
   case INBOUNDS:
static std::string toString(AssumptionKind Kind) {
   case DELINEARIZATION:
     return "Delinearization";
   case ERROR_DOMAINCONJUNCTS:
     return "Low number of domain conjuncts";
   }
   llvm_unreachable("Unknown AssumptionKind!");
 }
 
-void Scop::trackAssumption(AssumptionKind Kind, __isl_keep isl_set *Set,
-                           DebugLoc Loc) {
+bool Scop::trackAssumption(AssumptionKind Kind, __isl_keep isl_set *Set,
+                           DebugLoc Loc, bool IsPos) {
+  if (IsPos) {
     if (isl_set_is_subset(Context, Set))
-    return;
+      return false;
 
     if (isl_set_is_subset(AssumedContext, Set))
-    return;
+      return false;
+  } else {
+    if (isl_set_is_disjoint(Set, Context))
+      return false;
+
+    if (isl_set_is_subset(Set, InvalidContext))
+      return false;
+  }
 
   auto &F = *getRegion().getEntry()->getParent();
-  std::string Msg = toString(Kind) + " assumption:\t" + stringFromIslObj(Set);
+  auto Suffix = IsPos ? " assumption:\t" : " restriction:\t";
+  std::string Msg = toString(Kind) + Suffix + stringFromIslObj(Set);
   emitOptimizationRemarkAnalysis(F.getContext(), DEBUG_TYPE, F, Loc, Msg);
+  return true;
 }
 
 void Scop::addAssumption(AssumptionKind Kind, __isl_take isl_set *Set,
-                         DebugLoc Loc) {
-  trackAssumption(Kind, Set, Loc);
-  AssumedContext = isl_set_intersect(AssumedContext, Set);
-
-  int NSets = isl_set_n_basic_set(AssumedContext);
-  if (NSets >= MaxDisjunctsAssumed) {
-    isl_space *Space = isl_set_get_space(AssumedContext);
-    isl_set_free(AssumedContext);
-    AssumedContext = isl_set_empty(Space);
+                         DebugLoc Loc, bool IsPos) {
+  if (!trackAssumption(Kind, Set, Loc, IsPos)) {
+    isl_set_free(Set);
+    return;
   }
 
+  if (IsPos) {
+    AssumedContext = isl_set_intersect(AssumedContext, Set);
     AssumedContext = isl_set_coalesce(AssumedContext);
+  } else {
+    InvalidContext = isl_set_union(InvalidContext, Set);
+    InvalidContext = isl_set_coalesce(InvalidContext);
+  }
 }
 
 void Scop::invalidate(AssumptionKind Kind, DebugLoc Loc) {
-  addAssumption(Kind, isl_set_empty(getParamSpace()), Loc);
+  addAssumption(Kind, isl_set_empty(getParamSpace()), Loc, true);
 }
 
-__isl_give isl_set *Scop::getBoundaryContext() const {
-  return isl_set_copy(BoundaryContext);
+__isl_give isl_set *Scop::getInvalidContext() const {
+  return isl_set_copy(InvalidContext);
 }
 
 void Scop::printContext(raw_ostream &OS) const {
   OS << "Context:\n";
-
-  if (!Context) {
-    OS.indent(4) << "n/a\n\n";
-    return;
-  }
-
-  OS.indent(4) << getContextStr() << "\n";
+  OS.indent(4) << Context << "\n";
 
   OS.indent(4) << "Assumed Context:\n";
-  if (!AssumedContext) {
-    OS.indent(4) << "n/a\n\n";
-    return;
-  }
-
-  OS.indent(4) << getAssumedContextStr() << "\n";
-
-  OS.indent(4) << "Boundary Context:\n";
-  if (!BoundaryContext) {
-    OS.indent(4) << "n/a\n\n";
-    return;
-  }
+  OS.indent(4) << AssumedContext << "\n";
 
-  OS.indent(4) << getBoundaryContextStr() << "\n";
+  OS.indent(4) << "Invalid Context:\n";
+  OS.indent(4) << InvalidContext << "\n";
 
   for (const SCEV *Parameter : Parameters) {
     int Dim = ParameterIds.find(Parameter)->second;
     OS.indent(4) << "p" << Dim << ": " << *Parameter << "\n";
   }
 }
 
 void Scop::printAliasAssumptions(raw_ostream &OS) const {

lib/CodeGen/IslAst.cpp

buildCondition(__isl_keep isl_ast_build *Build, const Scop::MinMaxAccessTy *It0,
 
   return NonAliasGroup;
 }
 
 void IslAst::buildRunCondition(__isl_keep isl_ast_build *Build) {
   // The conditions that need to be checked at run-time for this scop are
   // available as an isl_set in the runtime check context from which we can
   // directly derive a run-time condition.
-  RunCondition =
-      isl_ast_build_expr_from_set(Build, S->getRuntimeCheckContext());
+  auto *OneV = isl_val_one(isl_ast_build_get_ctx(Build));
+  auto *PosCond = isl_ast_build_expr_from_set(Build, S->getAssumedContext());
+  auto *NegCond = isl_ast_build_expr_from_set(Build, S->getInvalidContext());
+  auto *NotNegCond = isl_ast_expr_sub(isl_ast_expr_from_val(OneV), NegCond);
+  RunCondition = isl_ast_expr_and(PosCond, NotNegCond);
 
   // Create the alias checks from the minimal/maximal accesses in each alias
   // group which consists of read only and non read only (read write) accesses.
   // This operation is by construction quadratic in the read-write pointers and
   // linear int the read only pointers in each alias group.
   for (const Scop::MinMaxVectorPairTy &MinMaxAccessPair : S->getAliasGroups()) {
     auto &MinMaxReadWrite = MinMaxAccessPair.first;
     auto &MinMaxReadOnly = MinMaxAccessPair.second;

test/Isl/Ast/OpenMP/nested_loop_both_parallel_parametric.ll

loop.i.backedge:
   %i.next = add nsw i64 %i, 1
   br label %loop.i
 
 ret:
   fence seq_cst
   ret void
 }
 
-; CHECK: if (n <= 1024)
+; CHECK: if (n <= 1024 && 1 - 0)
 ; CHECK:   #pragma omp parallel for
 ; CHECK:   for (int c0 = 0; c0 < n; c0 += 1)
 ; CHECK:     #pragma simd
 ; CHECK:     for (int c1 = 0; c1 < n; c1 += 1)
 ; CHECK:       Stmt_loop_body(c0, c1);

test/Isl/Ast/alias_simple_1.ll

 ; RUN: opt %loadPolly -polly-ast -analyze -disable-basicaa < %s | FileCheck %s --check-prefix=NOAA
 ; RUN: opt %loadPolly -polly-ast -analyze < %s | FileCheck %s --check-prefix=BASI
 ; RUN: opt %loadPolly -polly-ast -analyze -disable-basicaa -tbaa < %s | FileCheck %s --check-prefix=TBAA
 ; RUN: opt %loadPolly -polly-ast -analyze -disable-basicaa -scev-aa < %s | FileCheck %s --check-prefix=SCEV
 ; RUN: opt %loadPolly -polly-ast -analyze -disable-basicaa -globals-aa < %s | FileCheck %s --check-prefix=GLOB
 ;
 ;    int A[1024];
 ;
 ;
 ;    void jd(float *B, int N) {
 ;      for (int i = 0; i < N; i++)
 ;        A[i] = B[i];
 ;    }
 ;
-; NOAA: if (1 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
-; BASI: if (1 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
-; TBAA: if (1)
-; SCEV: if (1 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
-; GLOB: if (1 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
+; NOAA: if (1 && 1 - 0 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
+; BASI: if (1 && 1 - 0 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
+; TBAA: if (1 && 1 - 0)
+; SCEV: if (1 && 1 - 0 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
+; GLOB: if (1 && 1 - 0 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
 ;
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 @A = common global [1024 x i32] zeroinitializer, align 16
 
 define void @jd(float* nocapture readonly %B, i32 %N) {
 entry:
   %cmp6 = icmp sgt i32 %N, 0

test/Isl/Ast/alias_simple_2.ll

 ; RUN: opt %loadPolly -polly-ast -analyze -disable-basicaa < %s | FileCheck %s --check-prefix=NOAA
 ; RUN: opt %loadPolly -polly-ast -analyze < %s | FileCheck %s --check-prefix=BASI
 ; RUN: opt %loadPolly -polly-ast -analyze -disable-basicaa -tbaa < %s | FileCheck %s --check-prefix=TBAA
 ; RUN: opt %loadPolly -polly-ast -analyze -disable-basicaa -scev-aa < %s | FileCheck %s --check-prefix=SCEV
 ; RUN: opt %loadPolly -polly-ast -analyze -disable-basicaa -globals-aa < %s | FileCheck %s --check-prefix=GLOB
 ;
 ;    int A[1024], B[1024];
 ;
 ;
 ;    void jd(int N) {
 ;      for (int i = 0; i < N; i++)
 ;        A[i] = B[i];
 ;    }
 ;
-; NOAA: if (1 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
-; BASI: if (1)
-; TBAA: if (1 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
-; SCEV: if (1 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
-; GLOB: if (1 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
+; NOAA: if (1 && 1 - 0 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
+; BASI: if (1 && 1 - 0)
+; TBAA: if (1 && 1 - 0 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
+; SCEV: if (1 && 1 - 0 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
+; GLOB: if (1 && 1 - 0 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
 ;
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 @A = common global [1024 x i32] zeroinitializer, align 16
 @B = common global [1024 x i32] zeroinitializer, align 16
 
 define void @jd(i32 %N) {
 entry:

test/Isl/Ast/alias_simple_3.ll

 ; RUN: opt %loadPolly -polly-ast -analyze -disable-basicaa < %s | FileCheck %s --check-prefix=NOAA
 ; RUN: opt %loadPolly -polly-ast -analyze < %s | FileCheck %s --check-prefix=BASI
 ; RUN: opt %loadPolly -polly-ast -analyze -disable-basicaa -tbaa < %s | FileCheck %s --check-prefix=TBAA
 ; RUN: opt %loadPolly -polly-ast -analyze -disable-basicaa -scev-aa < %s | FileCheck %s --check-prefix=SCEV
 ; RUN: opt %loadPolly -polly-ast -analyze -disable-basicaa -globals-aa < %s | FileCheck %s --check-prefix=GLOB
 ;
 ;    int A[1024];
 ;    float B[1024];
 ;
 ;    void jd(int N) {
 ;      for (int i = 0; i < N; i++)
 ;        A[i] = B[i];
 ;    }
 ;
-; NOAA: if (1 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
-; BASI: if (1)
-; TBAA: if (1)
-; SCEV: if (1 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
-; GLOB: if (1 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
+; NOAA: if (1 && 1 - 0 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
+; BASI: if (1 && 1 - 0)
+; TBAA: if (1 && 1 - 0)
+; SCEV: if (1 && 1 - 0 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
+; GLOB: if (1 && 1 - 0 && (&MemRef_B[N] <= &MemRef_A[0] || &MemRef_A[N] <= &MemRef_B[0]))
 ;
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 @A = common global [1024 x i32] zeroinitializer, align 16
 @B = common global [1024 x float] zeroinitializer, align 16
 
 define void @jd(i32 %N) {
 entry:

test/Isl/Ast/aliasing_multiple_alias_groups.ll

 ; RUN: opt %loadPolly -polly-ast -analyze          < %s | FileCheck %s --check-prefix=NOAA
 ; RUN: opt %loadPolly -polly-ast -analyze -tbaa    < %s | FileCheck %s --check-prefix=TBAA
 ;
 ;    void jd(int *Int0, int *Int1, float *Float0, float *Float1) {
 ;      for (int i = 0; i < 1024; i++) {
 ;        Int0[i] = Int1[i];
 ;        Float0[i] = Float1[i];
 ;      }
 ;    }
 ;
-; NOAA:      if (1 && (
+; NOAA:      if (1 && 1 - 0 && (
 ; NOAA-DAG:    &MemRef_Int0[1024] <= &MemRef_Float0[0] || &MemRef_Float0[1024] <= &MemRef_Int0[0]
 ; NOAA-DAG:    &MemRef_Int1[1024] <= &MemRef_Float0[0] || &MemRef_Float0[1024] <= &MemRef_Int1[0]
 ; NOAA-DAG:    &MemRef_Float1[1024] <= &MemRef_Float0[0] || &MemRef_Float0[1024] <= &MemRef_Float1[0]
 ; NOAA-DAG:    &MemRef_Int1[1024] <= &MemRef_Int0[0] || &MemRef_Int0[1024] <= &MemRef_Int1[0]
 ; NOAA-DAG:    &MemRef_Float1[1024] <= &MemRef_Int0[0] || &MemRef_Int0[1024] <= &MemRef_Float1[0]
 ; NOAA:      ))
 ;
-; TBAA:      if (1 && (
+; TBAA:      if (1 && 1 - 0 && (
 ; TBAA-DAG:    &MemRef_Int1[1024] <= &MemRef_Int0[0] || &MemRef_Int0[1024] <= &MemRef_Int1[0]
 ; TBAA-DAG:    &MemRef_Float1[1024] <= &MemRef_Float0[0] || &MemRef_Float0[1024] <= &MemRef_Float1[0]
 ; TBAA:      ))
 ;
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 define void @jd(i32* nocapture %Int0, i32* nocapture readonly %Int1, float* nocapture %Float0, float* nocapture readonly %Float1) {
 entry:

test/Isl/Ast/aliasing_parametric_simple_1.ll

 ; RUN: opt %loadPolly -polly-ast -analyze %s | FileCheck %s
 ;
 ;    void jd(int *A, int *B, int c) {
 ;      for (int i = 0; i < 1024; i++)
 ;        A[i] = B[c];
 ;    }
 ;
-; CHECK: if (1 && (&MemRef_B[c + 1] <= &MemRef_A[0] || &MemRef_A[1024] <= &MemRef_B[c]))
+; CHECK: if (1 && 1 - 0 && (&MemRef_B[c + 1] <= &MemRef_A[0] || &MemRef_A[1024] <= &MemRef_B[c]))
 ; CHECK:     for (int c0 = 0; c0 <= 1023; c0 += 1)
 ; CHECK:       Stmt_for_body(c0);
 ; CHECK: else
 ; CHECK:     /* original code */
 ;
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 define void @jd(i32* %A, i32* %B, i32 %c) {

test/Isl/Ast/aliasing_parametric_simple_2.ll

 ; RUN: opt %loadPolly -polly-ast -analyze < %s | FileCheck %s
 ;
 ;    void jd(int *A, int *B, int c) {
 ;      for (int i = 0; i < 1024; i++)
 ;        A[i] = B[c - 10] + B[5];
 ;    }
 ;
-; CHECK: if (c >= -{{[0-9]*}} && (&MemRef_B[c <= 15 ? 6 : c - 9] <= &MemRef_A[0] || &MemRef_A[1024] <= &MemRef_B[c >= 15 ? 5 : c - 10]))
+; CHECK: if (1 && 1 - (c <= -{{[0-9]*}}) && (&MemRef_B[c <= 15 ? 6 : c - 9] <= &MemRef_A[0] || &MemRef_A[1024] <= &MemRef_B[c >= 15 ? 5 : c - 10]))
 ; CHECK:     for (int c0 = 0; c0 <= 1023; c0 += 1)
 ; CHECK:       Stmt_for_body(c0);
 ; CHECK: else
 ; CHECK:    /* original code */
 ;
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 define void @jd(i32* %A, i32* %B, i32 %c) {

test/Isl/Ast/run-time-condition.ll

for.body:                                         ; preds = %for.cond
   store i16 %add10, i16* %arrayidx13, align 2
   %inc = add nsw i64 %indvar, 1
   br label %for.cond
 
 for.end:                                          ; preds = %for.cond
   ret void
 }
 
-; CHECK: if (1)
+; CHECK: if (1 && 1 - 0)
 ; CHECK:     for (int c0 = 0; c0 <= 1023; c0 += 1)
 ; CHECK:       Stmt_for_body(c0);
 ; CHECK: else
 ; CHECK:     {  /* original code */ }
 

test/Isl/Ast/simple-run-time-condition.ll

 ;     for (long j = 0; j < m; j++)
 ;         A[i+p][j+q-100] = 1.0;
 ;
 
 ; This test case is meant to verify that the run-time condition generated
 ; for the delinearization is simplified such that conditions that would not
 ; cause any code to be executed are not generated.
 
-; CHECK: if (
-; CHECK: (o >= 1 && n + p <= 9223372036854775808 && q <= 0 && m + q >= 0)
-; CHECK: ||
-; CHECK; (o <= 0 && m + q >= 100 && q <= 100)
-; CHECK: )
+; CHECK: if (((o >= 1 && q <= 0 && m + q >= 0) || (o <= 0 && m + q >= 100 && q <= 100)) && 1 - (n + p >= 9223372036854775809 || (o <= 0 && m + q >= 9223372036854775909) || (o <= 0 && q <= -9223372036854775709)))
 
 ; CHECK:     if (o <= 0) {
 ; CHECK:       for (int c0 = 0; c0 < n; c0 += 1)
 ; CHECK:         for (int c1 = 0; c1 < m; c1 += 1)
 ; CHECK:           Stmt_for_j_1(c0, c1);
 ; CHECK:     } else
 ; CHECK:       for (int c0 = 0; c0 < n; c0 += 1)
 ; CHECK:         for (int c1 = 0; c1 < m; c1 += 1)

test/Isl/CodeGen/aliasing_parametric_simple_2.ll

 ; RUN: opt %loadPolly -polly-codegen -S < %s | FileCheck %s
 ;
 ;    void jd(int *A, int *B, int c) {
 ;      for (int i = 0; i < 1024; i++)
 ;        A[i] = B[c - 10] + B[5];
 ;    }
 ;
 ; CHECK:  sext i32 %c to i64
 ; CHECK:  sext i32 %c to i64
+; CHECK:  %[[Ctx:[._a-zA-Z0-9]*]] = and i1 true
 ; CHECK:  %[[M0:[._a-zA-Z0-9]*]] = sext i32 %c to i64
 ; CHECK:  %[[M1:[._a-zA-Z0-9]*]] = icmp sle i64 %[[M0]], 15
 ; CHECK:  %[[M2:[._a-zA-Z0-9]*]] = sext i32 %c to i64
 ; CHECK:  %[[M3:[._a-zA-Z0-9]*]] = sub nsw i64 %[[M2]], 9
 ; CHECK:  %[[M4:[._a-zA-Z0-9]*]] = select i1 %[[M1]], i64 6, i64 %[[M3]]
 ; CHECK:  %[[BMax:[._a-zA-Z0-9]*]] = getelementptr i32, i32* %B, i64 %[[M4]]
 ; CHECK:  %[[AMin:[._a-zA-Z0-9]*]] = getelementptr i32, i32* %A, i64 0
 ; CHECK:  %[[BltA:[._a-zA-Z0-9]*]] = icmp ule i32* %[[BMax]], %[[AMin]]
 ; CHECK:  %[[AMax:[._a-zA-Z0-9]*]] = getelementptr i32, i32* %A, i64 1024
 ; CHECK:  %[[m0:[._a-zA-Z0-9]*]] = sext i32 %c to i64
 ; CHECK:  %[[m1:[._a-zA-Z0-9]*]] = icmp sge i64 %[[m0]], 15
 ; CHECK:  %[[m2:[._a-zA-Z0-9]*]] = sext i32 %c to i64
 ; CHECK:  %[[m3:[._a-zA-Z0-9]*]] = sub nsw i64 %[[m2]], 10
 ; CHECK:  %[[m4:[._a-zA-Z0-9]*]] = select i1 %[[m1]], i64 5, i64 %[[m3]]
 ; CHECK:  %[[BMin:[._a-zA-Z0-9]*]] = getelementptr i32, i32* %B, i64 %[[m4]]
 ; CHECK:  %[[AltB:[._a-zA-Z0-9]*]] = icmp ule i32* %[[AMax]], %[[BMin]]
 ; CHECK:  %[[NoAlias:[._a-zA-Z0-9]*]] = or i1 %[[BltA]], %[[AltB]]
-; CHECK:  %[[RTC:[._a-zA-Z0-9]*]] = and i1 %3, %[[NoAlias]]
+; CHECK:  %[[RTC:[._a-zA-Z0-9]*]] = and i1 %[[Ctx]], %[[NoAlias]]
 ; CHECK:  br i1 %[[RTC]], label %polly.start, label %for.cond
 ;
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 define void @jd(i32* %A, i32* %B, i32 %c) {
 entry:
   br label %for.cond
 

test/Isl/CodeGen/exprModDiv.ll

 
 ; #define floord(n,d) ((n < 0) ? (n - d + 1) : n) / d
 ; A[p + 127 * floord(-p - 1, 127) + 127]
 ; CHECK:  %pexp.fdiv_q.0 = sub i64 %p, 127
 ; CHECK:  %pexp.fdiv_q.1 = add i64 %pexp.fdiv_q.0, 1
 ; CHECK:  %pexp.fdiv_q.2 = icmp slt i64 %p, 0
 ; CHECK:  %pexp.fdiv_q.3 = select i1 %pexp.fdiv_q.2, i64 %pexp.fdiv_q.1, i64 %p
 ; CHECK:  %pexp.fdiv_q.4 = sdiv i64 %pexp.fdiv_q.3, 127
-; CHECK:  %17 = mul nsw i64 127, %pexp.fdiv_q.4
-; CHECK:  %18 = sub nsw i64 %p, %17
-; CHECK:  %polly.access.A8 = getelementptr float, float* %A, i64 %18
+; CHECK:  %[[r1:[0-9]*]] = mul nsw i64 127, %pexp.fdiv_q.4
+; CHECK:  %[[r2:[0-9]*]] = sub nsw i64 %p, %[[r1]]
+; CHECK:  %polly.access.A8 = getelementptr float, float* %A, i64 %[[r2]]
 
 ; A[p / 127]
 ; CHECK:  %pexp.div = sdiv exact i64 %p, 127
 ; CHECK:  %polly.access.B9 = getelementptr float, float* %B, i64 %pexp.div
 
 ; A[i % 128]
 ; POW2:  %pexp.pdiv_r = urem i64 %polly.indvar, 128
 ; POW2:  %polly.access.A6 = getelementptr float, float* %A, i64 %pexp.pdiv_r
 
 ; A[floor(i / 128)]
 ; POW2:  %pexp.p_div_q = udiv i64 %polly.indvar, 128
 ; POW2:  %polly.access.B7 = getelementptr float, float* %B, i64 %pexp.p_div_q
 
 ; #define floord(n,d) ((n < 0) ? (n - d + 1) : n) / d
 ; A[p + 128 * floord(-p - 1, 128) + 128]
 ; POW2:  %polly.fdiv_q.shr = ashr i64 %p, 7
-; POW2:  %17 = mul nsw i64 128, %polly.fdiv_q.shr
-; POW2:  %18 = sub nsw i64 %p, %17
-; POW2:  %polly.access.A8 = getelementptr float, float* %A, i64 %18
+; POW2:  %[[r1:[0-9]*]] = mul nsw i64 128, %polly.fdiv_q.shr
+; POW2:  %[[r2:[0-9]*]] = sub nsw i64 %p, %[[r1]]
+; POW2:  %polly.access.A8 = getelementptr float, float* %A, i64 %[[r2]]
 
 ; A[p / 128]
 ; POW2:  %pexp.div = sdiv exact i64 %p, 128
 ; POW2:  %polly.access.B9 = getelementptr float, float* %B, i64 %pexp.div
 
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 define void @exprModDiv(float* %A, float* %B, float* %C, i64 %N, i64 %p) {

test/Isl/CodeGen/large-numbers-in-boundary-context.ll

 ; RUN: opt %loadPolly -S -polly-codegen < %s | FileCheck %s
 ;
 ; The boundary context contains a constant that does not fit in 64 bits. Hence,
 ; we will check that we use an appropriaty typed constant, here with 65 bits.
 ; An alternative would be to bail out early but that would not be as easy.
 ;
-; CHECK: {{.*}} = icmp sge i65 {{.*}}, -9223372036854775809
+; CHECK: {{.*}} = icmp sle i65 {{.*}}, -9223372036854775810
 ;
 ; CHECK: polly.start
 ;
 target triple = "x86_64-unknown-linux-gnu"
 
 @global = external global i32, align 4
 @global1 = external global i32, align 4
 

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

 
 ; RUN: opt %loadPolly -polly-ast -analyze -polly-ignore-aliasing < %s | FileCheck %s -check-prefix=SCALAR
 ; RUN: opt %loadPolly -polly-codegen -polly-ignore-aliasing -disable-output < %s
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 ; SCALAR: if (
 ; SCALAR:     {
 ; 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: if (1 && 1 - 0)
 ; SCALAR:     Stmt_for_0(0);
 
 
 define void @foo(i32* %A) {
 entry:
   %Scalar0 = alloca i32
   br label %for.0
 

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

 ; SCALAR-NEXT: Region: %bb1---%bb13
 ; SCALAR-NEXT: Max Loop Depth:  1
 ; SCALAR-NEXT: Invariant Accesses: {
 ; SCALAR-NEXT: }
 ; SCALAR-NEXT: Context:
 ; SCALAR-NEXT: {  :  }
 ; SCALAR-NEXT: Assumed Context:
 ; SCALAR-NEXT: {  :  }
-; SCALAR-NEXT: Boundary Context:
-; SCALAR-NEXT: {  :  }
+; SCALAR-NEXT: Invalid Context:
+; SCALAR-NEXT: {  : 1 = 0 }
 ; SCALAR-NEXT: Arrays {
 ; SCALAR-NEXT:     i32 MemRef_C[*]; // Element size 4
 ; SCALAR-NEXT:     i32 MemRef_A[*]; // Element size 4
 ; SCALAR-NEXT: }
 ; SCALAR-NEXT: Arrays (Bounds as pw_affs) {
 ; SCALAR-NEXT:     i32 MemRef_C[*]; // Element size 4
 ; SCALAR-NEXT:     i32 MemRef_A[*]; // Element size 4
 ; SCALAR-NEXT: }

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

 ; INNERMOST-NEXT: Max Loop Depth:  1
 ; INNERMOST-NEXT: Invariant Accesses: {
 ; INNERMOST-NEXT: }
 ; INNERMOST-NEXT: Context:
 ; INNERMOST-NEXT: [p_0, p_1, p_2] -> {  : 0 <= p_0 <= 1048576 and 0 <= p_1 <= 1024 and 0 <= p_2 <= 1024 }
 ; INNERMOST-NEXT: Assumed Context:
 ; INNERMOST-NEXT: [p_0, p_1, p_2] -> {  :  }
 ; INNERMOST-NEXT: Boundary Context:
-; INNERMOST-NEXT: [p_0, p_1, p_2] -> {  :  }
+; INNERMOST-NEXT: [p_0, p_1, p_2] -> {  : 1 = 0 }
 ; INNERMOST-NEXT: p0: {0,+,{0,+,1}<nuw><nsw><%bb11>}<nuw><nsw><%bb13>
 ; INNERMOST-NEXT: p1: {0,+,1}<nuw><nsw><%bb11>
 ; INNERMOST-NEXT: p2: {0,+,1}<nuw><nsw><%bb13>
 ; INNERMOST-NEXT: Arrays {
 ; INNERMOST-NEXT:     i32 MemRef_A[*]; // Element size 4
 ; INNERMOST-NEXT:     i64 MemRef_indvars_iv_next6; // Element size 8
 ; INNERMOST-NEXT:     i64 MemRef_indvars_iv_next4; // Element size 8
 ; INNERMOST-NEXT: }

 ; ALL-NEXT: Region: %bb11---%bb29
 ; ALL-NEXT: Max Loop Depth:  2
 ; ALL-NEXT: Invariant Accesses: {
 ; ALL-NEXT: }
 ; ALL-NEXT: Context:
 ; ALL-NEXT: {  :  }
 ; ALL-NEXT: Assumed Context:
 ; ALL-NEXT: {  :  }
-; ALL-NEXT: Boundary Context:
-; ALL-NEXT: {  :  }
+; ALL-NEXT: Invalid Context:
+; ALL-NEXT: {  : 1 = 0 }
 ; ALL-NEXT: Arrays {
 ; ALL-NEXT:     i32 MemRef_A[*]; // Element size 4
 ; ALL-NEXT: }
 ; ALL-NEXT: Arrays (Bounds as pw_affs) {
 ; ALL-NEXT:     i32 MemRef_A[*]; // Element size 4
 ; ALL-NEXT: }
 ; ALL-NEXT: Alias Groups (0):
 ; ALL-NEXT:     n/a

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

 ; INNERMOST-NEXT: Region: %bb15---%bb13
 ; INNERMOST-NEXT: Max Loop Depth:  1
 ; INNERMOST-NEXT: Invariant Accesses: {
 ; INNERMOST-NEXT: }
 ; INNERMOST-NEXT: Context:
 ; INNERMOST-NEXT: [p_0, p_1, p_2] -> {  : 0 <= p_0 <= 2147483647 and 0 <= p_1 <= 1024 and 0 <= p_2 <= 1024 }
 ; INNERMOST-NEXT: Assumed Context:
 ; INNERMOST-NEXT: [p_0, p_1, p_2] -> {  :  }
-; INNERMOST-NEXT: Boundary Context:
-; INNERMOST-NEXT: [p_0, p_1, p_2] -> {  :  }
+; INNERMOST-NEXT: Invalid Context:
+; INNERMOST-NEXT: [p_0, p_1, p_2] -> {  : 1 = 0 }
 ; INNERMOST-NEXT: p0: {0,+,{0,+,1}<nuw><nsw><%bb11>}<nuw><nsw><%bb13>
 ; INNERMOST-NEXT: p1: {0,+,1}<nuw><nsw><%bb11>
 ; INNERMOST-NEXT: p2: {0,+,1}<nuw><nsw><%bb13>
 ; INNERMOST-NEXT: Arrays {
 ; INNERMOST-NEXT:     i32 MemRef_A[*]; // Element size 4
 ; INNERMOST-NEXT:     i64 MemRef_indvars_iv_next6; // Element size 8
 ; INNERMOST-NEXT:     i32 MemRef_indvars_iv_next4; // Element size 4
 ; INNERMOST-NEXT: }

 ; ALL-NEXT: Region: %bb11---%bb29
 ; ALL-NEXT: Max Loop Depth:  2
 ; ALL-NEXT: Invariant Accesses: {
 ; ALL-NEXT: }
 ; ALL-NEXT: Context:
 ; ALL-NEXT: {  :  }
 ; ALL-NEXT: Assumed Context:
 ; ALL-NEXT: {  :  }
-; ALL-NEXT: Boundary Context:
-; ALL-NEXT: {  :  }
+; ALL-NEXT: Invalid Context:
+; ALL-NEXT: {  : 1 = 0 }
 ; ALL-NEXT: Arrays {
 ; ALL-NEXT:     i32 MemRef_A[*]; // Element size 4
 ; ALL-NEXT: }
 ; ALL-NEXT: Arrays (Bounds as pw_affs) {
 ; ALL-NEXT:     i32 MemRef_A[*]; // Element size 4
 ; ALL-NEXT: }
 ; ALL-NEXT: Alias Groups (0):
 ; ALL-NEXT:     n/a

test/ScopInfo/NonAffine/non_affine_conditional_surrounding_affine_loop.ll

 ; ALL-NEXT: Region: %bb3---%bb19
 ; ALL-NEXT: Max Loop Depth:  1
 ; ALL-NEXT: Invariant Accesses: {
 ; ALL-NEXT: }
 ; ALL-NEXT: Context:
 ; ALL-NEXT: {  :  }
 ; ALL-NEXT: Assumed Context:
 ; ALL-NEXT: {  :  }
-; ALL-NEXT: Boundary Context:
-; ALL-NEXT: {  :  }
+; ALL-NEXT: Invalid Context:
+; ALL-NEXT: {  : 1 = 0 }
 ; ALL-NEXT: Arrays {
 ; ALL-NEXT:     i32 MemRef_A[*]; // Element size 4
 ; ALL-NEXT: }
 ; ALL-NEXT: Arrays (Bounds as pw_affs) {
 ; ALL-NEXT:     i32 MemRef_A[*]; // Element size 4
 ; ALL-NEXT: }
 ; ALL-NEXT: Alias Groups (0):
 ; ALL-NEXT:     n/a

test/ScopInfo/NonAffine/non_affine_conditional_surrounding_non_affine_loop.ll

 ; ALL-NEXT: Region: %bb3---%bb20
 ; ALL-NEXT: Max Loop Depth:  1
 ; ALL-NEXT: Invariant Accesses: {
 ; ALL-NEXT: }
 ; ALL-NEXT: Context:
 ; ALL-NEXT: {  :  }
 ; ALL-NEXT: Assumed Context:
 ; ALL-NEXT: {  :  }
-; ALL-NEXT: Boundary Context:
-; ALL-NEXT: {  :  }
+; ALL-NEXT: Invalid Context:
+; ALL-NEXT: {  : 1 = 0 }
 ; ALL-NEXT: Arrays {
 ; ALL-NEXT:     i32 MemRef_A[*]; // Element size 4
 ; ALL-NEXT: }
 ; ALL-NEXT: Arrays (Bounds as pw_affs) {
 ; ALL-NEXT:     i32 MemRef_A[*]; // Element size 4
 ; ALL-NEXT: }
 ; ALL-NEXT: Alias Groups (0):
 ; ALL-NEXT:     n/a

test/ScopInfo/NonAffine/non_affine_float_compare.ll

 ; CHECK-NEXT: Region: %bb1---%bb14
 ; CHECK-NEXT: Max Loop Depth:  1
 ; CHECK-NEXT: Invariant Accesses: {
 ; CHECK-NEXT: }
 ; CHECK-NEXT: Context:
 ; CHECK-NEXT: {  :  }
 ; CHECK-NEXT: Assumed Context:
 ; CHECK-NEXT: {  :  }
-; CHECK-NEXT: Boundary Context:
-; CHECK-NEXT: {  :  }
+; CHECK-NEXT: Invalid Context:
+; CHECK-NEXT: {  : 1 = 0 }
 ; CHECK-NEXT: Arrays {
 ; CHECK-NEXT:     float MemRef_A[*]; // Element size 4
 ; CHECK-NEXT: }
 ; CHECK-NEXT: Arrays (Bounds as pw_affs) {
 ; CHECK-NEXT:     float MemRef_A[*]; // Element size 4
 ; CHECK-NEXT: }
 ; CHECK-NEXT: Alias Groups (0):
 ; CHECK-NEXT:     n/a

test/ScopInfo/NonAffine/non_affine_loop_condition.ll

 ; CHECK-NEXT: Region: %bb1---%bb12
 ; CHECK-NEXT: Max Loop Depth:  1
 ; CHECK-NEXT: Invariant Accesses: {
 ; CHECK-NEXT: }
 ; CHECK-NEXT: Context:
 ; CHECK-NEXT: {  :  }
 ; CHECK-NEXT: Assumed Context:
 ; CHECK-NEXT: {  :  }
-; CHECK-NEXT: Boundary Context:
-; CHECK-NEXT: {  :  }
+; CHECK-NEXT: Invalid Context:
+; CHECK-NEXT: {  : 1 = 0 }
 ; CHECK-NEXT: Arrays {
 ; CHECK-NEXT:     i32 MemRef_C[*]; // Element size 4
 ; CHECK-NEXT:     i32 MemRef_A[*]; // Element size 4
 ; CHECK-NEXT: }
 ; CHECK-NEXT: Arrays (Bounds as pw_affs) {
 ; CHECK-NEXT:     i32 MemRef_C[*]; // Element size 4
 ; CHECK-NEXT:     i32 MemRef_A[*]; // Element size 4
 ; CHECK-NEXT: }

test/ScopInfo/NonAffine/non_affine_loop_used_later.ll

 ; CHECK-NEXT: Region: %bb2---%bb24
 ; CHECK-NEXT: Max Loop Depth:  1
 ; CHECK-NEXT: Invariant Accesses: {
 ; CHECK-NEXT: }
 ; CHECK-NEXT: Context:
 ; CHECK-NEXT: [N] -> {  : -2147483648 <= N <= 2147483647 }
 ; CHECK-NEXT: Assumed Context:
 ; CHECK-NEXT: [N] -> {  :  }
-; CHECK-NEXT: Boundary Context:
-; CHECK-NEXT: [N] -> {  :  }
+; CHECK-NEXT: Invalid Context:
+; CHECK-NEXT: [N] -> {  : 1 = 0 }
 ; CHECK-NEXT: p0: %N
 ; CHECK-NEXT: Arrays {
 ; CHECK-NEXT:     i32 MemRef_j_0__phi; // Element size 4
 ; CHECK-NEXT:     i32 MemRef_j_0; // Element size 4
 ; CHECK-NEXT:     i32 MemRef_A[*]; // Element size 4
 ; CHECK-NEXT:     i32 MemRef_j_2__phi; // Element size 4
 ; CHECK-NEXT:     i32 MemRef_j_2; // Element size 4
 ; CHECK-NEXT: }

test/ScopInfo/assume_gep_bounds_2.ll

 ;    }
 
 ; This code is within bounds either if m and p are smaller than the array sizes,
 ; but also if only p is smaller than the size of the second B dimension and n
 ; is such that the first loop is never executed and consequently A is never
 ; accessed. In this case the value of m does not matter.
 
 ; CHECK:      Assumed Context:
-; CHECK-NEXT: [n, m, p] -> {  : p <= 20 and (n <= 0 or (n > 0 and m <= 20)) }
+; CHECK-NEXT:   [n, m, p] -> { : p <= 20 and (n <= 0 or (n > 0 and m <= 20)) }
 
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 define void @foo([20 x float]* noalias %A, [20 x float]* noalias %B, i64 %n, i64 %m, i64 %p) {
 entry:
   br label %for.cond
 
 for.cond:                                         ; preds = %for.inc5, %entry

test/ScopInfo/bounded_loop_assumptions.ll

 ; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
 ;
 ; The assumed context is tricky here as the equality test for the inner loop
 ; allows an "unbounded" loop trip count. We assume that does not happen, thus
 ; if N <= 1 the outer loop is not executed and we are done, if N >= 3 the
 ; equality test in the inner exit condition will trigger at some point and,
 ; finally, if N == 2 we would have an unbounded inner loop.
 ;
 ; CHECK:      Assumed Context:
-; CHECK-NEXT:   [N] -> {  : N >= 3 or N <= 1 }
+; CHECK-NEXT:   [N] -> {  :  }
+; CHECK:      Invalid Context:
+; CHECK-NEXT:   [N] -> {  : N = 2 }
 ;
 ;    int jd(int *restrict A, int x, int N) {
 ;      for (int i = 1; i < N; i++)
 ;        for (int j = 3; j < N; j++)
 ;          x += A[i];
 ;      return x;
 ;    }
 ;

test/ScopInfo/long-sequence-of-error-blocks-2.ll

 ; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
 
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "x86_64-unknown-linux-gnu"
 
 %struct.hoge = type { i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, [8 x [2 x i32]], [8 x [2 x i32]], [4 x [4 x i32]], i32, i32, i32, i32, [256 x i8], [256 x i8], [256 x i8], [256 x i8], [256 x i8], i32, i32, i32, i32, i32, i32, [500 x i8], i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, [256 x i8], [256 x i8], [256 x i8], i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, [1024 x i8], i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, double, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, [256 x i8], [256 x i8], i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, [256 x i8], i32, i32, i32*, i32*, i8*, i32*, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, double, double, double, [5 x double], i32, [8 x i32], i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, [6 x double], [6 x double], [256 x i8], i32, i32, i32, i32, [2 x [5 x i32]], [2 x [5 x i32]], i32, i32, i32, i32, i32, i32, i32, i32, i32, [3 x i32], i32 }
 
 ; The assumed context of this test case has at some point become very complex.
-; This test case verifies that we bail out and generate an infeasible
-; assumed context. It should be updated after having made the assumed context
-; construction more efficient.
-
-; CHECK-NOT: Assumed Context:
+; However, since we keep both the assumed as well as invalid context that
+; problem is solved.
+;
+; CHECK:      Assumed Context:
+; CHECK-NEXT:   [tmp17, tmp21, tmp27, tmp31, tmp37, tmp41, tmp46, tmp52, tmp56, tmp62] -> {  :  }
+; CHECK:      Invalid Context:
+; CHECK-NEXT:   [tmp17, tmp21, tmp27, tmp31, tmp37, tmp41, tmp46, tmp52, tmp56, tmp62] -> {  : (tmp17 < 0 and tmp21 < 0) or (tmp17 < 0 and tmp21 > 0) or (tmp17 > 0 and tmp21 < 0) or (tmp17 > 0 and tmp21 > 0) or (tmp37 < 0 and tmp41 < 0 and tmp46 > 0) or (tmp37 < 0 and tmp41 > 0 and tmp46 > 0) or (tmp37 > 0 and tmp41 < 0 and tmp46 > 0) or (tmp37 > 0 and tmp41 > 0 and tmp46 > 0) or (tmp27 = 3 and tmp31 <= 143) or (tmp56 = 0 and tmp52 < 0) or (tmp56 = 0 and tmp52 > 0) }
 
 
 @global = external global [300 x i8], align 16
 @global1 = external global %struct.hoge*, align 8
 @global2 = external unnamed_addr constant [79 x i8], align 1
 @global3 = external unnamed_addr constant [57 x i8], align 1
 
 declare void @widget() #0

test/ScopInfo/long-sequence-of-error-blocks.ll

 ; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
 
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "x86_64-unknown-linux-gnu"
 
 %struct.hoge = type { i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, [8 x [2 x i32]], [8 x [2 x i32]], [4 x [4 x i32]], i32, i32, i32, i32, [256 x i8], [256 x i8], [256 x i8], [256 x i8], [256 x i8], i32, i32, i32, i32, i32, i32, [500 x i8], i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, [256 x i8], [256 x i8], [256 x i8], i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, [1024 x i8], i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, double, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, [256 x i8], [256 x i8], i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, [256 x i8], i32, i32, i32*, i32*, i8*, i32*, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, double, double, double, [5 x double], i32, [8 x i32], i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, [6 x double], [6 x double], [256 x i8], i32, i32, i32, i32, [2 x [5 x i32]], [2 x [5 x i32]], i32, i32, i32, i32, i32, i32, i32, i32, i32, [3 x i32], i32 }
 
-; The assumed context of this test case is still very large and should be
-; reduced.
-
 ; This test case contains a long sequence of branch instructions together with
 ; function calls that are considered 'error blocks'. We verify that the
 ; iteration spaces are not overly complicated.
 
 ; CHECK:      Assumed Context:
-; CHECK-NEXT: [tmp17, tmp21, tmp27, tmp31] -> {  : (tmp17 = 0 and tmp27 >= 4) or (tmp17 = 0 and tmp27 <= 2) or (tmp21 = 0 and tmp17 < 0 and tmp27 >= 4) or (tmp21 = 0 and tmp17 < 0 and tmp27 <= 2) or (tmp21 = 0 and tmp17 > 0 and tmp27 >= 4) or (tmp21 = 0 and tmp17 > 0 and tmp27 <= 2) or (tmp17 = 0 and tmp27 = 3 and tmp31 >= 144) or (tmp21 = 0 and tmp27 = 3 and tmp17 < 0 and tmp31 >= 144) or (tmp21 = 0 and tmp27 = 3 and tmp17 > 0 and tmp31 >= 144) }
+; CHECK-NEXT: [tmp17, tmp21, tmp27, tmp31] -> {  :  }
+; CHECK:      Invalid Context:
+; CHECK-NEXT: [tmp17, tmp21, tmp27, tmp31] -> {  : (tmp17 < 0 and tmp21 < 0) or (tmp17 < 0 and tmp21 > 0) or (tmp17 > 0 and tmp21 < 0) or (tmp17 > 0 and tmp21 > 0) or (tmp27 = 3 and tmp31 <= 143) }
 ;
 ; CHECK:      Statements {
 ; CHECK-NEXT:     Stmt_bb15
 ; CHECK-NEXT:         Domain :=
 ; CHECK-NEXT:             [tmp17, tmp21, tmp27, tmp31] -> { Stmt_bb15[] };
 ; CHECK-NEXT:         Schedule :=
 ; CHECK-NEXT:             [tmp17, tmp21, tmp27, tmp31] -> { Stmt_bb15[] -> [0] };
 ; CHECK-NEXT:         MustWriteAccess :=    [Reduction Type: NONE] [Scalar: 0]

test/ScopInfo/multidim_3d_parametric_array_static_loop_bounds.ll

 ; RUN: opt %loadPolly -polly-scops -analyze -polly-delinearize < %s | FileCheck %s
 target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64-S128"
 
 ; void foo(long n, long m, long o, double A[n][m][o]) {
 ;
 ;   for (long i = 0; i < 100; i++)
 ;     for (long j = 0; j < 150; j++)
 ;       for (long k = 0; k < 200; k++)
 ;         A[i][j][k] = 1.0;
 ; }
 
 ; CHECK:      Assumed Context:
 ; CHECK-NEXT: [m, o] -> {  : m >= 150 and o >= 200 }
+; CHECK:      Invalid Context:
+; CHECK-NEXT: [m, o] -> {  : 1 = 0 }
 ;
 ; CHECK:      p0: %m
 ; CHECK-NEXT: p1: %o
 ;
 ; CHECK:      Statements {
 ; CHECK-NEXT:     Stmt_for_k
 ; CHECK-NEXT:         Domain :=
 ; CHECK-NEXT:             [m, o] -> { Stmt_for_k[i0, i1, i2] : 0 <= i0 <= 99 and 0 <= i1 <= 149 and 0 <= i2 <= 199 };

test/ScopInfo/non-pure-function-call.ll

 ; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
 ;
 ; CHECK:      Assumed Context:
-; CHECK-NEXT:   [N] -> {  : N <= 101 }
+; CHECK-NEXT:   [N] -> {  :  }
+; CHECK:      Invalid Context:
+; CHECK-NEXT:   [N] -> {  : N >= 102 }
 ;
 ;    void g(void);
 ;    void f(int *A, int N) {
 ;      for (int i = 0; i < N; i++) {
 ;        if (i > 100)
 ;          g();
 ;        A[i]++;
 ;      }

test/ScopInfo/non-pure-function-calls-causes-dead-blocks.ll

 ;        timer_stop();
 ;    }
 ;
 ;  The assumed context should not be empty even though all statements are
 ;  executed only if timeit != 0.
 ;
 ; CHECK:    Region: %entry.split---%if.end.20
 ; CHECK:    Assumed Context:
-; CHECK-NEXT:    [timeit, N] -> { : timeit = 0 }
+; CHECK-NEXT:    [timeit, N] -> { : }
+; CHECK:    Invalid Context:
+; CHECK-NEXT:    [timeit, N] -> { : timeit < 0 or timeit > 0 }
 ; CHECK:    Statements {
 ; CHECK-NOT:      Stmt_if_then_split
 ; CHECK:      Stmt_for_body
 ; CHECK:      Stmt_for_body_9
 ; CHECK:    }
 ;
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 

test/ScopInfo/non-pure-function-calls.ll

 ;
 ;      for (int i = 0; i < N; i++)
 ;        A[i] += B[i];
 ;
 ;      if (timeit)
 ;        timer_stop();
 ;    }
 ;
-; CHECK:    Region: %for.cond---%if.end.20
-; CHECK:    Assumed Context:
-; CHECK:    [N, timeit] -> {  : timeit = 0 }
-; CHECK:    Statements {
-; CHECK:      Stmt
-; CHECK:      Stmt
-; CHECK-NOT   Stmt
-; CHECK:    }
+; CHECK:      Region: %for.cond---%if.end.20
+; CHECK:      Assumed Context:
+; CHECK-NEXT: [N, timeit] -> {  :  }
+; CHECK:      Invalid Context:
+; CHECK-NEXT: [N, timeit] -> {  : timeit < 0 or timeit > 0 }
+; CHECK:      Statements {
+; CHECK:        Stmt
+; CHECK:        Stmt
+; CHECK-NOT:    Stmt
+; CHECK:      }
 ;
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 define void @kernel(i32* %A, i32* %B, i32 %timeit, i32 %N) {
 entry:
   %tobool = icmp eq i32 %timeit, 0
   br i1 %tobool, label %if.end, label %if.then
 

test/ScopInfo/remarks.ll

 ; RUN: opt %loadPolly -pass-remarks-analysis="polly-scops" -polly-scops -disable-output < %s 2>&1 | FileCheck %s
 ;
 ; CHECK: remark: test/ScopInfo/remarks.c:4:7: SCoP begins here.
-; CHECK: remark: test/ScopInfo/remarks.c:8:5: Finite loop assumption:    [M, N, Debug] -> {  : M >= 0 or (M <= -2 and N <= 0) or (M = -1 and N <= 0) }
-; CHECK: remark: test/ScopInfo/remarks.c:13:7: No-error assumption:    [M, N, Debug] -> {  : M < 0 or (M >= 0 and N <= 0) or (Debug = 0 and M >= 0 and N > 0) }
-; CHECK: remark: test/ScopInfo/remarks.c:9:7: Inbounds assumption:    [M, N, Debug] -> {  : M <= 100 or (M > 0 and N <= 0) }
-; CHECK: remark: <unknown>:0:0: No-overflows assumption:    [N, M, Debug] -> {  : -2147483648 - N <= M <= 2147483647 - N }
+; CHECK: remark: test/ScopInfo/remarks.c:8:5: Finite loop restriction:  [M, N, Debug] -> {  : N > 0 and (M <= -2 or M = -1) }
+; CHECK: remark: test/ScopInfo/remarks.c:13:7: No-error restriction:  [M, N, Debug] -> {  : M >= 0 and N > 0 and (Debug < 0 or Debug > 0) }
+; CHECK: remark: test/ScopInfo/remarks.c:9:7: Inbounds assumption:  [M, N, Debug] -> {  : M <= 100 or (M > 0 and N <= 0) }
+; CHECK: remark: <unknown>:0:0: No-overflows restriction:  [N, M, Debug] -> {  : M <= -2147483649 - N or M >= 2147483648 - N }
 ; CHECK: remark: test/ScopInfo/remarks.c:9:18: Possibly aliasing pointer, use restrict keyword.
 ; CHECK: remark: test/ScopInfo/remarks.c:9:33: Possibly aliasing pointer, use restrict keyword.
 ; CHECK: remark: test/ScopInfo/remarks.c:9:15: Possibly aliasing pointer, use restrict keyword.
 ; CHECK: remark: test/ScopInfo/remarks.c:14:3: SCoP ends here.
 ; CHECK: remark: test/ScopInfo/remarks.c:19:3: SCoP begins here.
 ; CHECK: remark: test/ScopInfo/remarks.c:21:11: Invariant load assumption:    [tmp] -> {  : 1 = 0 }
 ; CHECK: remark: test/ScopInfo/remarks.c:22:16: SCoP ends here but was dismissed.
 ;

test/ScopInfo/switch-1.ll

 ; CHECK-NEXT:         Schedule :=
 ; CHECK-NEXT:             [N] -> { Stmt_sw_bb_6[i0] -> [i0, 0] };
 ; CHECK-NEXT:         ReadAccess :=    [Reduction Type: +] [Scalar: 0]
 ; CHECK-NEXT:             [N] -> { Stmt_sw_bb_6[i0] -> MemRef_A[i0] };
 ; CHECK-NEXT:         MustWriteAccess :=    [Reduction Type: +] [Scalar: 0]
 ; CHECK-NEXT:             [N] -> { Stmt_sw_bb_6[i0] -> MemRef_A[i0] };
 ; CHECK-NEXT: }
 
-; AST:      if (1)
+; AST:      if (1 && 1 - 0)
 ;
 ; AST:          {
 ; AST-NEXT:       for (int c0 = 1; c0 < N - 2; c0 += 4) {
 ; AST-NEXT:         Stmt_sw_bb_1(c0);
 ; AST-NEXT:         Stmt_sw_bb_2(c0 + 1);
 ; AST-NEXT:         Stmt_sw_bb_6(c0 + 2);
 ; AST-NEXT:       }
 ; AST-NEXT:       if (N >= 2)

test/ScopInfo/switch-2.ll

 ; CHECK-NEXT:         Schedule :=
 ; CHECK-NEXT:             [N] -> { Stmt_sw_bb_2[i0] -> [i0, 0] };
 ; CHECK-NEXT:         ReadAccess :=    [Reduction Type: +] [Scalar: 0]
 ; CHECK-NEXT:             [N] -> { Stmt_sw_bb_2[i0] -> MemRef_A[i0] };
 ; CHECK-NEXT:         MustWriteAccess :=    [Reduction Type: +] [Scalar: 0]
 ; CHECK-NEXT:             [N] -> { Stmt_sw_bb_2[i0] -> MemRef_A[i0] };
 ; CHECK-NEXT: }
 
-; AST:      if (1)
+; AST:      if (1 && 1 - 0)
 ;
 ; AST:          for (int c0 = 0; c0 < N; c0 += 4) {
 ; AST-NEXT:       Stmt_sw_bb(c0);
 ; AST-NEXT:       if (N >= c0 + 3)
 ; AST-NEXT:         Stmt_sw_bb_2(c0 + 2);
 ; AST-NEXT:     }
 ;
 ; AST:      else

test/ScopInfo/switch-3.ll

 ; CHECK-NEXT:         Schedule :=
 ; CHECK-NEXT:             [N] -> { Stmt_sw_bb_9[i0] -> [i0, 1] };
 ; CHECK-NEXT:         ReadAccess :=    [Reduction Type: +] [Scalar: 0]
 ; CHECK-NEXT:             [N] -> { Stmt_sw_bb_9[i0] -> MemRef_A[i0] };
 ; CHECK-NEXT:         MustWriteAccess :=    [Reduction Type: +] [Scalar: 0]
 ; CHECK-NEXT:             [N] -> { Stmt_sw_bb_9[i0] -> MemRef_A[i0] };
 ; CHECK-NEXT: }
 
-; AST:      if (1)
+; AST:      if (1 && 1 - 0)
 ;
 ; AST:          for (int c0 = 0; c0 < N; c0 += 1) {
 ; AST-NEXT:       if ((c0 - 2) % 4 == 0)
 ; AST-NEXT:         Stmt_sw_bb_5(c0);
 ; AST-NEXT:       if (c0 % 4 >= 2) {
 ; AST-NEXT:         Stmt_sw_bb_9(c0);
 ; AST-NEXT:       } else {
 ; AST-NEXT:         if (c0 % 4 == 0)

test/ScopInfo/switch-4.ll

 ; CHECK-NEXT:         Schedule :=
 ; CHECK-NEXT:             [N] -> { Stmt_sw_bb_9[i0] -> [i0, 0] };
 ; CHECK-NEXT:         ReadAccess :=    [Reduction Type: +] [Scalar: 0]
 ; CHECK-NEXT:             [N] -> { Stmt_sw_bb_9[i0] -> MemRef_A[i0] };
 ; CHECK-NEXT:         MustWriteAccess :=    [Reduction Type: +] [Scalar: 0]
 ; CHECK-NEXT:             [N] -> { Stmt_sw_bb_9[i0] -> MemRef_A[i0] };
 ; CHECK-NEXT: }
 
-; AST:      if (1)
+; AST:      if (1 && 1 - 0)
 ;
 ; AST:          {
 ; AST-NEXT:       for (int c0 = 0; c0 < N - 3; c0 += 4) {
 ; AST-NEXT:         Stmt_sw_bb(c0);
 ; AST-NEXT:         Stmt_sw_bb_1(c0 + 1);
 ; AST-NEXT:         Stmt_sw_bb_5(c0 + 2);
 ; AST-NEXT:         Stmt_sw_bb_9(c0 + 3);
 ; AST-NEXT:       }

test/ScopInfo/switch-5.ll

 ; CHECK-NEXT:         Schedule :=
 ; CHECK-NEXT:             [N] -> { Stmt_for_body[i0] -> [0] };
 ; CHECK-NEXT:         ReadAccess :=    [Reduction Type: +] [Scalar: 0]
 ; CHECK-NEXT:             [N] -> { Stmt_for_body[i0] -> MemRef_A[0] };
 ; CHECK-NEXT:         MustWriteAccess :=    [Reduction Type: +] [Scalar: 0]
 ; CHECK-NEXT:             [N] -> { Stmt_for_body[i0] -> MemRef_A[0] };
 ; CHECK-NEXT: }
 
-; AST:      if (1)
+; AST:      if (1 && 1 - 0)
 ;
 ; AST:          if (N >= 1)
 ; AST-NEXT:       Stmt_for_body(0);
 ;
 ; AST:      else
 ; AST-NEXT:     {  /* original code */ }
 
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"

test/ScopInfo/switch-6.ll

 ; CHECK-NEXT:         Schedule :=
 ; CHECK-NEXT:             [N] -> { Stmt_sw_bb_9[i0] -> [3, 0] };
 ; CHECK-NEXT:         ReadAccess :=    [Reduction Type: +] [Scalar: 0]
 ; CHECK-NEXT:             [N] -> { Stmt_sw_bb_9[i0] -> MemRef_A[3] };
 ; CHECK-NEXT:         MustWriteAccess :=    [Reduction Type: +] [Scalar: 0]
 ; CHECK-NEXT:             [N] -> { Stmt_sw_bb_9[i0] -> MemRef_A[3] };
 ; CHECK-NEXT: }
 
-; AST:      if (1)
+; AST:      if (1 && 1 - 0)
 ;
 ; AST:          if (N >= 1) {
 ; AST-NEXT:       Stmt_sw_bb(0);
 ; AST-NEXT:       if (N >= 2) {
 ; AST-NEXT:         Stmt_sw_bb_1(1);
 ; AST-NEXT:         if (N >= 3) {
 ; AST-NEXT:           Stmt_sw_bb_5(2);
 ; AST-NEXT:           if (N >= 4)

test/ScopInfo/switch-7.ll

 ; CHECK-NEXT:         ReadAccess :=    [Reduction Type: NONE] [Scalar: 0]
 ; CHECK-NEXT:             [c, N] -> { Stmt_for_body_7[i0] -> MemRef_A[i0] };
 ; CHECK-NEXT:         ReadAccess :=    [Reduction Type: NONE] [Scalar: 0]
 ; CHECK-NEXT:             [c, N] -> { Stmt_for_body_7[i0] -> MemRef_A[1 + i0] };
 ; CHECK-NEXT:         MustWriteAccess :=    [Reduction Type: NONE] [Scalar: 0]
 ; CHECK-NEXT:             [c, N] -> { Stmt_for_body_7[i0] -> MemRef_A[1 + i0] };
 ; CHECK-NEXT: }
 
-; AST:      if (1)
+; AST:      if (1 && 1 - 0)
 ;
 ; AST:          if (c == 1) {
 ; AST-NEXT:       for (int c0 = 0; c0 < N; c0 += 1)
 ; AST-NEXT:         Stmt_for_body_7(c0);
 ; AST-NEXT:     } else if (c == -1)
 ; AST-NEXT:       for (int c0 = 0; c0 < N; c0 += 1)
 ; AST-NEXT:         Stmt_for_body(c0);
 ;

test/ScopInfo/test-wrapping-in-condition.ll

 ; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
 ;
-; CHECK:    Boundary Context:
-; CHECK:        [N] -> {  : N <= 128 }
+; CHECK:    Invalid Context:
+; CHECK:        [N] -> {  : N >= 129 }
 ;
 ;    #include <stdlib.h>
 ;    #include <stdio.h>
 ;
 ;    void __attribute__((noinline)) foo(float *A, long N) {
 ;      for (long i = 0; i < N; i++)
 ;        if ((signed char)i < 100)
 ;          A[i] += i;

test/ScopInfo/user_provided_assumptions.ll

 ; RUN: opt %loadPolly -pass-remarks-analysis="polly-scops" -polly-scops -disable-output < %s 2>&1 | FileCheck %s
 ; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s --check-prefix=SCOP
 ;
 ; CHECK:      remark: <unknown>:0:0: SCoP begins here.
 ; CHECK-NEXT: remark: <unknown>:0:0: Use user assumption: [M, N] -> {  : N <= 2147483647 - M }
 ; CHECK-NEXT: remark: <unknown>:0:0: Use user assumption: [M, N] -> {  : -2147483648 - M <= N <= 2147483647 - M }
 ; CHECK-NEXT: remark: <unknown>:0:0: Use user assumption: [M, N, Debug] -> {  : Debug = 0 and 0 < M <= 100 and -2147483648 - M <= N <= 2147483647 - M }
 ; CHECK-NEXT: remark: <unknown>:0:0: Use user assumption: [M, N, Debug] -> {  : Debug = 0 and 0 < M <= 100 and N > 0 and -2147483648 - M <= N <= 2147483647 - M }
 ; CHECK-NEXT: remark: <unknown>:0:0: SCoP ends here.
 
 ; SCOP:      Context:
 ; SCOP-NEXT: [N, M, Debug] -> {  : Debug = 0 and N > 0 and 0 < M <= 2147483647 - N and M <= 100 }
 ; SCOP:      Assumed Context:
 ; SCOP-NEXT: [N, M, Debug] -> {  :  }
-; SCOP:      Boundary Context:
-; SCOP-NEXT: [N, M, Debug] -> {  :  }
+; SCOP:      Invalid Context:
+; SCOP-NEXT: [N, M, Debug] -> {  : 1 = 0 }
 ;
 ;    #include <stdio.h>
 ;
 ;    void valid(int * restrict A, int * restrict B, int N, int M, int C[100][100], int Debug) {
 ;      __builtin_assume(M <= 2147483647 - N);
 ;      __builtin_assume(M >= -2147483648 - N);
 ;      __builtin_assume(Debug == 0 && M <= 100 && M >= 1 && N >= 1);
 ;      if (N + M == -1)

test/ScopInfo/wraping_signed_expr_0.ll

 ; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
 ;
 ;    void f(int *A, char N, char p) {
 ;      for (char i = 0; i < N; i++) {
 ;        A[i + 3] = 0;
 ;      }
 ;    }
 ;
 ; The wrap function has no inbounds GEP but the nowrap function has. Therefore,
 ; we will add the assumption that i+1 won't overflow only to the former.
 ;
 ; CHECK:      Function: wrap
-; CHECK:      Boundary Context:
-; CHECK:      [N] -> {  : N <= 125 }
+; CHECK:      Invalid Context:
+; CHECK:      [N] -> {  : N >= 126 }
 ;
 ;
 ; FIXME: This is a negative test as nowrap should not need an assumed context.
 ;        However %tmp5 in @nowrap is translated to the SCEV <3,+,1><nw><%bb2>
 ;        which lacks the <nsw> flags we would need to avoid runtime checks.
 ;
 ; CHECK:      Function: nowrap
-; CHECK:      Boundary Context:
+; CHECK:      Invalid Context:
 ; CHECK-NOT:  [N] -> {  :  }
 ;
 target datalayout = "e-m:e-i8:64-f80:128-n8:16:32:64-S128"
 
 define void @wrap(i32* %A, i8 %N, i8 %p) {
 bb:
   br label %bb2
 

test/ScopInfo/wraping_signed_expr_1.ll

 ; Note:
 ;       1152921504606846975 * sizeof(long) <= 2 ^ 63 - 1
 ;  and
 ;       1152921504606846976 * sizeof(long) >  2 ^ 63 - 1
 ; with
 ;       sizeof(long) == 8
 ;
 ; CHECK:      Function: wrap
-; CHECK:      Boundary Context:
-; CHECK:      [N] -> {  : N <= 1152921504606846975 }
+; CHECK:      Invalid Context:
+; CHECK-NEXT: [N] -> {  : N >= 1152921504606846976 }
 ;
 ; CHECK:      Function: nowrap
-; CHECK:      Boundary Context:
-; CHECK:      [N] -> {  :  }
+; CHECK:      Invalid Context:
+; CHECK-NEXT: [N] -> {  : 1 = 0 }
 ;
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 define void @wrap(i64* %A, i64 %N, i64 %p) {
 bb:
   br label %bb2
 
 bb2:                                              ; preds = %bb7, %bb

test/ScopInfo/wraping_signed_expr_2.ll

 ;
 ; Note: 2147483618 + 30 == 2 ^ 31
 
 ; CHECK:      Function: wrap
 ;
 ; CHECK:      Context:
 ; CHECK-NEXT: [N] -> {  : -2147483648 <= N <= 2147483647 }
 ;
-; CHECK:      Boundary Context:
-; CHECK-NEXT: [N] -> {  : N <= 2147483618 }
+; CHECK:      Invalid Context:
+; CHECK-NEXT: [N] -> {  : N >= 2147483619 }
 
 target datalayout = "e-m:e-i32:64-f80:128-n8:16:32:64-S128"
 
 define void @wrap(i32* %A, i32 %N, i32 %p) {
 bb:
   br label %bb2
 
 bb2:                                              ; preds = %bb7, %bb

test/ScopInfo/wraping_signed_expr_3.ll

 ; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
 ;
 ;    void f(int *A, int N, int p) {
 ;      for (int i = 0; i < N; i++)
 ;        A[i + p] = 0;
 ;    }
 ;
 ; Note: 2147483648 == 2 ^ 31
 ;
 ; CHECK:      Function: wrap
-; CHECK:      Boundary Context:
-; CHECK:      [N, p] -> {  : p <= 2147483648 - N }
+; CHECK:      Invalid Context:
+; CHECK:      [N, p] -> {  : p >= 2147483649 - N }
 ;
 target datalayout = "e-m:e-i32:64-f80:128-n8:16:32:64-S128"
 
 define void @wrap(i32* %A, i32 %N, i32 %p) {
 bb:
   br label %bb2
 
 bb2:                                              ; preds = %bb7, %bb

test/ScopInfo/wraping_signed_expr_4.ll

 ; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
 ;
 ;    void f(char *A, char N, char p) {
 ;      for (char i = 0; i < N; i++)
 ;        A[p-1] = 0;
 ;    }
 
 ; CHECK:      Function: wrap
 ;
 ; CHECK:      Context:
 ; CHECK-NEXT: [N, p] -> {  : -128 <= N <= 127 and -128 <= p <= 127 }
 ;
-; CHECK:      Boundary Context:
-; CHECK-NEXT: [N, p] -> {  : p >= -127 }
+; CHECK:      Invalid Context:
+; CHECK-NEXT: [N, p] -> {  : p = -128 }
 
 target datalayout = "e-m:e-i8:64-f80:128-n8:16:32:64-S128"
 
 define void @wrap(i8* %A, i8 %N, i8 %p) {
 bb:
   br label %bb2
 
 bb2:                                              ; preds = %bb7, %bb

test/ScopInfo/wraping_signed_expr_5.ll

 ; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
 ;
 ; We should not generate runtime check for ((int)r1 + (int)r2) as it is known not
 ; to overflow. However (p + q) can, thus checks are needed.
 ;
-; CHECK:      Boundary Context:
-; CHECK-NEXT: [r1, r2, q, p] -> {  : r2 <= 127 + r1 and -2147483648 - q <= p <= 2147483647 - q }
+; CHECK:      Invalid Context:
+; CHECK-NEXT: [r1, r2, q, p] -> { : p <= -2147483649 - q or p >= 2147483648 - q or r2 >= 128 + 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"
 

test/ScopInfo/wraping_signed_expr_6.ll

 ; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
 ;
-; CHECK:    Boundary Context:
-; CHECK:        [N] -> {  : N <= 128 }
+; CHECK:    Invalid Context:
+; CHECK:        [N] -> {  : N >= 129 }
 ;
 ;    void foo(float *A, long N) {
 ;      for (long i = 0; i < N; i++)
 ;        if ((signed char)i < 100)
 ;          A[i] += i;
 ;    }
 define void @foo(float* %A, i64 %N) {
 bb:

test/ScopInfo/wraping_signed_expr_7.ll

 ; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
 ;
-; CHECK:    Boundary Context:
-; CHECK:        [N] -> {  : N <= 128 }
+; CHECK:    Invalid Context:
+; CHECK:        [N] -> {  : N >= 129 }
 ;
 ;    void foo(float *A, long N) {
 ;      for (long i = 0; i < N;)
 ;        if ((signed char)i++ < 100)
 ;          A[i] += i;
 ;    }
 define void @foo(float* %A, i64 %N) {
 bb: