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