polly/lib/CodeGen/IslAst.cpp

   1 //===- IslAst.cpp - isl code generator interface --------------------------===//
   2 //
   3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
   4 // See https://llvm.org/LICENSE.txt for license information.
   5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
   6 //
   7 //===----------------------------------------------------------------------===//
   8 //
   9 // The isl code generator interface takes a Scop and generates an isl_ast. This
  10 // ist_ast can either be returned directly or it can be pretty printed to
  11 // stdout.
  12 //
  13 // A typical isl_ast output looks like this:
  14 //
  15 // for (c2 = max(0, ceild(n + m, 2); c2 <= min(511, floord(5 * n, 3)); c2++) {
  16 //   bb2(c2);
  17 // }
  18 //
  19 // An in-depth discussion of our AST generation approach can be found in:
  20 //
  21 // Polyhedral AST generation is more than scanning polyhedra
  22 // Tobias Grosser, Sven Verdoolaege, Albert Cohen
  23 // ACM Transactions on Programming Languages and Systems (TOPLAS),
  24 // 37(4), July 2015
  25 // http://www.grosser.es/#pub-polyhedral-AST-generation
  26 //
  27 //===----------------------------------------------------------------------===//
  28
  29 #include "polly/CodeGen/IslAst.h"
  30 #include "polly/CodeGen/CodeGeneration.h"
  31 #include "polly/DependenceInfo.h"
  32 #include "polly/LinkAllPasses.h"
  33 #include "polly/Options.h"
  34 #include "polly/ScopDetection.h"
  35 #include "polly/ScopInfo.h"
  36 #include "polly/ScopPass.h"
  37 #include "polly/Support/GICHelper.h"
  38 #include "llvm/ADT/Statistic.h"
  39 #include "llvm/IR/Function.h"
  40 #include "llvm/Support/Debug.h"
  41 #include "llvm/Support/raw_ostream.h"
  42 #include "isl/aff.h"
  43 #include "isl/ast.h"
  44 #include "isl/ast_build.h"
  45 #include "isl/id.h"
  46 #include "isl/isl-noexceptions.h"
  47 #include "isl/printer.h"
  48 #include "isl/schedule.h"
  49 #include "isl/set.h"
  50 #include "isl/union_map.h"
  51 #include "isl/val.h"
  52 #include <cassert>
  53 #include <cstdlib>
  54
  55 #define DEBUG_TYPE "polly-ast"
  56
  57 using namespace llvm;
  58 using namespace polly;
  59
  60 using IslAstUserPayload = IslAstInfo::IslAstUserPayload;
  61
  62 static cl::opt<bool>
  63     PollyParallel("polly-parallel",
  64                   cl::desc("Generate thread parallel code (isl codegen only)"),
  65                   cl::cat(PollyCategory));
  66
  67 static cl::opt<bool> PrintAccesses("polly-ast-print-accesses",
  68                                    cl::desc("Print memory access functions"),
  69                                    cl::cat(PollyCategory));
  70
  71 static cl::opt<bool> PollyParallelForce(
  72     "polly-parallel-force",
  73     cl::desc(
  74         "Force generation of thread parallel code ignoring any cost model"),
  75     cl::cat(PollyCategory));
  76
  77 static cl::opt<bool> UseContext("polly-ast-use-context",
  78                                 cl::desc("Use context"), cl::Hidden,
  79                                 cl::init(true), cl::cat(PollyCategory));
  80
  81 static cl::opt<bool> DetectParallel("polly-ast-detect-parallel",
  82                                     cl::desc("Detect parallelism"), cl::Hidden,
  83                                     cl::cat(PollyCategory));
  84
  85 STATISTIC(ScopsProcessed, "Number of SCoPs processed");
  86 STATISTIC(ScopsBeneficial, "Number of beneficial SCoPs");
  87 STATISTIC(BeneficialAffineLoops, "Number of beneficial affine loops");
  88 STATISTIC(BeneficialBoxedLoops, "Number of beneficial boxed loops");
  89
  90 STATISTIC(NumForLoops, "Number of for-loops");
  91 STATISTIC(NumParallel, "Number of parallel for-loops");
  92 STATISTIC(NumInnermostParallel, "Number of innermost parallel for-loops");
  93 STATISTIC(NumOutermostParallel, "Number of outermost parallel for-loops");
  94 STATISTIC(NumReductionParallel, "Number of reduction-parallel for-loops");
  95 STATISTIC(NumExecutedInParallel, "Number of for-loops executed in parallel");
  96 STATISTIC(NumIfConditions, "Number of if-conditions");
  97
  98 namespace polly {
  99
 100 /// Temporary information used when building the ast.
 101 struct AstBuildUserInfo {
 102   /// Construct and initialize the helper struct for AST creation.
 103   AstBuildUserInfo() = default;
 104
 105   /// The dependence information used for the parallelism check.
 106   const Dependences *Deps = nullptr;
 107
 108   /// Flag to indicate that we are inside a parallel for node.
 109   bool InParallelFor = false;
 110
 111   /// Flag to indicate that we are inside an SIMD node.
 112   bool InSIMD = false;
 113
 114   /// The last iterator id created for the current SCoP.
 115   isl_id *LastForNodeId = nullptr;
 116 };
 117 } // namespace polly
 118
 119 /// Free an IslAstUserPayload object pointed to by @p Ptr.
 120 static void freeIslAstUserPayload(void *Ptr) {
 121   delete ((IslAstInfo::IslAstUserPayload *)Ptr);
 122 }
 123
 124 /// Print a string @p str in a single line using @p Printer.
 125 static isl_printer *printLine(__isl_take isl_printer *Printer,
 126                               const std::string &str,
 127                               __isl_keep isl_pw_aff *PWA = nullptr) {
 128   Printer = isl_printer_start_line(Printer);
 129   Printer = isl_printer_print_str(Printer, str.c_str());
 130   if (PWA)
 131     Printer = isl_printer_print_pw_aff(Printer, PWA);
 132   return isl_printer_end_line(Printer);
 133 }
 134
 135 /// Return all broken reductions as a string of clauses (OpenMP style).
 136 static const std::string getBrokenReductionsStr(const isl::ast_node &Node) {
 137   IslAstInfo::MemoryAccessSet *BrokenReductions;
 138   std::string str;
 139
 140   BrokenReductions = IslAstInfo::getBrokenReductions(Node);
 141   if (!BrokenReductions || BrokenReductions->empty())
 142     return "";
 143
 144   // Map each type of reduction to a comma separated list of the base addresses.
 145   std::map<MemoryAccess::ReductionType, std::string> Clauses;
 146   for (MemoryAccess *MA : *BrokenReductions)
 147     if (MA->isWrite())
 148       Clauses[MA->getReductionType()] +=
 149           ", " + MA->getScopArrayInfo()->getName();
 150
 151   // Now print the reductions sorted by type. Each type will cause a clause
 152   // like:  reduction (+ : sum0, sum1, sum2)
 153   for (const auto &ReductionClause : Clauses) {
 154     str += " reduction (";
 155     str += MemoryAccess::getReductionOperatorStr(ReductionClause.first);
 156     // Remove the first two symbols (", ") to make the output look pretty.
 157     str += " : " + ReductionClause.second.substr(2) + ")";
 158   }
 159
 160   return str;
 161 }
 162
 163 /// Callback executed for each for node in the ast in order to print it.
 164 static isl_printer *cbPrintFor(__isl_take isl_printer *Printer,
 165                                __isl_take isl_ast_print_options *Options,
 166                                __isl_keep isl_ast_node *Node, void *) {
 167   isl::pw_aff DD =
 168       IslAstInfo::getMinimalDependenceDistance(isl::manage_copy(Node));
 169   const std::string BrokenReductionsStr =
 170       getBrokenReductionsStr(isl::manage_copy(Node));
 171   const std::string KnownParallelStr = "#pragma known-parallel";
 172   const std::string DepDisPragmaStr = "#pragma minimal dependence distance: ";
 173   const std::string SimdPragmaStr = "#pragma simd";
 174   const std::string OmpPragmaStr = "#pragma omp parallel for";
 175
 176   if (!DD.is_null())
 177     Printer = printLine(Printer, DepDisPragmaStr, DD.get());
 178
 179   if (IslAstInfo::isInnermostParallel(isl::manage_copy(Node)))
 180     Printer = printLine(Printer, SimdPragmaStr + BrokenReductionsStr);
 181
 182   if (IslAstInfo::isExecutedInParallel(isl::manage_copy(Node)))
 183     Printer = printLine(Printer, OmpPragmaStr);
 184   else if (IslAstInfo::isOutermostParallel(isl::manage_copy(Node)))
 185     Printer = printLine(Printer, KnownParallelStr + BrokenReductionsStr);
 186
 187   return isl_ast_node_for_print(Node, Printer, Options);
 188 }
 189
 190 /// Check if the current scheduling dimension is parallel.
 191 ///
 192 /// In case the dimension is parallel we also check if any reduction
 193 /// dependences is broken when we exploit this parallelism. If so,
 194 /// @p IsReductionParallel will be set to true. The reduction dependences we use
 195 /// to check are actually the union of the transitive closure of the initial
 196 /// reduction dependences together with their reversal. Even though these
 197 /// dependences connect all iterations with each other (thus they are cyclic)
 198 /// we can perform the parallelism check as we are only interested in a zero
 199 /// (or non-zero) dependence distance on the dimension in question.
 200 static bool astScheduleDimIsParallel(const isl::ast_build &Build,
 201                                      const Dependences *D,
 202                                      IslAstUserPayload *NodeInfo) {
 203   if (!D->hasValidDependences())
 204     return false;
 205
 206   isl::union_map Schedule = Build.get_schedule();
 207   isl::union_map Dep = D->getDependences(
 208       Dependences::TYPE_RAW | Dependences::TYPE_WAW | Dependences::TYPE_WAR);
 209
 210   if (!D->isParallel(Schedule.get(), Dep.release())) {
 211     isl::union_map DepsAll =
 212         D->getDependences(Dependences::TYPE_RAW | Dependences::TYPE_WAW |
 213                           Dependences::TYPE_WAR | Dependences::TYPE_TC_RED);
 214     // TODO: We will need to change isParallel to stop the unwrapping
 215     isl_pw_aff *MinimalDependenceDistanceIsl = nullptr;
 216     D->isParallel(Schedule.get(), DepsAll.release(),
 217                   &MinimalDependenceDistanceIsl);
 218     NodeInfo->MinimalDependenceDistance =
 219         isl::manage(MinimalDependenceDistanceIsl);
 220     return false;
 221   }
 222
 223   isl::union_map RedDeps = D->getDependences(Dependences::TYPE_TC_RED);
 224   if (!D->isParallel(Schedule.get(), RedDeps.release()))
 225     NodeInfo->IsReductionParallel = true;
 226
 227   if (!NodeInfo->IsReductionParallel)
 228     return true;
 229
 230   for (const auto &MaRedPair : D->getReductionDependences()) {
 231     if (!MaRedPair.second)
 232       continue;
 233     isl::union_map MaRedDeps = isl::manage_copy(MaRedPair.second);
 234     if (!D->isParallel(Schedule.get(), MaRedDeps.release()))
 235       NodeInfo->BrokenReductions.insert(MaRedPair.first);
 236   }
 237   return true;
 238 }
 239
 240 // This method is executed before the construction of a for node. It creates
 241 // an isl_id that is used to annotate the subsequently generated ast for nodes.
 242 //
 243 // In this function we also run the following analyses:
 244 //
 245 // - Detection of openmp parallel loops
 246 //
 247 static __isl_give isl_id *astBuildBeforeFor(__isl_keep isl_ast_build *Build,
 248                                             void *User) {
 249   AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
 250   IslAstUserPayload *Payload = new IslAstUserPayload();
 251   isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", Payload);
 252   Id = isl_id_set_free_user(Id, freeIslAstUserPayload);
 253   BuildInfo->LastForNodeId = Id;
 254
 255   Payload->IsParallel = astScheduleDimIsParallel(isl::manage_copy(Build),
 256                                                  BuildInfo->Deps, Payload);
 257
 258   // Test for parallelism only if we are not already inside a parallel loop
 259   if (!BuildInfo->InParallelFor && !BuildInfo->InSIMD)
 260     BuildInfo->InParallelFor = Payload->IsOutermostParallel =
 261         Payload->IsParallel;
 262
 263   return Id;
 264 }
 265
 266 // This method is executed after the construction of a for node.
 267 //
 268 // It performs the following actions:
 269 //
 270 // - Reset the 'InParallelFor' flag, as soon as we leave a for node,
 271 //   that is marked as openmp parallel.
 272 //
 273 static __isl_give isl_ast_node *
 274 astBuildAfterFor(__isl_take isl_ast_node *Node, __isl_keep isl_ast_build *Build,
 275                  void *User) {
 276   isl_id *Id = isl_ast_node_get_annotation(Node);
 277   assert(Id && "Post order visit assumes annotated for nodes");
 278   IslAstUserPayload *Payload = (IslAstUserPayload *)isl_id_get_user(Id);
 279   assert(Payload && "Post order visit assumes annotated for nodes");
 280
 281   AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
 282   assert(Payload->Build.is_null() && "Build environment already set");
 283   Payload->Build = isl::manage_copy(Build);
 284   Payload->IsInnermost = (Id == BuildInfo->LastForNodeId);
 285
 286   Payload->IsInnermostParallel =
 287       Payload->IsInnermost && (BuildInfo->InSIMD || Payload->IsParallel);
 288   if (Payload->IsOutermostParallel)
 289     BuildInfo->InParallelFor = false;
 290
 291   isl_id_free(Id);
 292   return Node;
 293 }
 294
 295 static isl_stat astBuildBeforeMark(__isl_keep isl_id *MarkId,
 296                                    __isl_keep isl_ast_build *Build,
 297                                    void *User) {
 298   if (!MarkId)
 299     return isl_stat_error;
 300
 301   AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
 302   if (strcmp(isl_id_get_name(MarkId), "SIMD") == 0)
 303     BuildInfo->InSIMD = true;
 304
 305   return isl_stat_ok;
 306 }
 307
 308 static __isl_give isl_ast_node *
 309 astBuildAfterMark(__isl_take isl_ast_node *Node,
 310                   __isl_keep isl_ast_build *Build, void *User) {
 311   assert(isl_ast_node_get_type(Node) == isl_ast_node_mark);
 312   AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
 313   auto *Id = isl_ast_node_mark_get_id(Node);
 314   if (strcmp(isl_id_get_name(Id), "SIMD") == 0)
 315     BuildInfo->InSIMD = false;
 316   isl_id_free(Id);
 317   return Node;
 318 }
 319
 320 static __isl_give isl_ast_node *AtEachDomain(__isl_take isl_ast_node *Node,
 321                                              __isl_keep isl_ast_build *Build,
 322                                              void *User) {
 323   assert(!isl_ast_node_get_annotation(Node) && "Node already annotated");
 324
 325   IslAstUserPayload *Payload = new IslAstUserPayload();
 326   isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", Payload);
 327   Id = isl_id_set_free_user(Id, freeIslAstUserPayload);
 328
 329   Payload->Build = isl::manage_copy(Build);
 330
 331   return isl_ast_node_set_annotation(Node, Id);
 332 }
 333
 334 // Build alias check condition given a pair of minimal/maximal access.
 335 static isl::ast_expr buildCondition(Scop &S, isl::ast_build Build,
 336                                     const Scop::MinMaxAccessTy *It0,
 337                                     const Scop::MinMaxAccessTy *It1) {
 338
 339   isl::pw_multi_aff AFirst = It0->first;
 340   isl::pw_multi_aff ASecond = It0->second;
 341   isl::pw_multi_aff BFirst = It1->first;
 342   isl::pw_multi_aff BSecond = It1->second;
 343
 344   isl::id Left = AFirst.get_tuple_id(isl::dim::set);
 345   isl::id Right = BFirst.get_tuple_id(isl::dim::set);
 346
 347   isl::ast_expr True =
 348       isl::ast_expr::from_val(isl::val::int_from_ui(Build.ctx(), 1));
 349   isl::ast_expr False =
 350       isl::ast_expr::from_val(isl::val::int_from_ui(Build.ctx(), 0));
 351
 352   const ScopArrayInfo *BaseLeft =
 353       ScopArrayInfo::getFromId(Left)->getBasePtrOriginSAI();
 354   const ScopArrayInfo *BaseRight =
 355       ScopArrayInfo::getFromId(Right)->getBasePtrOriginSAI();
 356   if (BaseLeft && BaseLeft == BaseRight)
 357     return True;
 358
 359   isl::set Params = S.getContext();
 360
 361   isl::ast_expr NonAliasGroup, MinExpr, MaxExpr;
 362
 363   // In the following, we first check if any accesses will be empty under
 364   // the execution context of the scop and do not code generate them if this
 365   // is the case as isl will fail to derive valid AST expressions for such
 366   // accesses.
 367
 368   if (!AFirst.intersect_params(Params).domain().is_empty() &&
 369       !BSecond.intersect_params(Params).domain().is_empty()) {
 370     MinExpr = Build.access_from(AFirst).address_of();
 371     MaxExpr = Build.access_from(BSecond).address_of();
 372     NonAliasGroup = MaxExpr.le(MinExpr);
 373   }
 374
 375   if (!BFirst.intersect_params(Params).domain().is_empty() &&
 376       !ASecond.intersect_params(Params).domain().is_empty()) {
 377     MinExpr = Build.access_from(BFirst).address_of();
 378     MaxExpr = Build.access_from(ASecond).address_of();
 379
 380     isl::ast_expr Result = MaxExpr.le(MinExpr);
 381     if (!NonAliasGroup.is_null())
 382       NonAliasGroup = isl::manage(
 383           isl_ast_expr_or(NonAliasGroup.release(), Result.release()));
 384     else
 385       NonAliasGroup = Result;
 386   }
 387
 388   if (NonAliasGroup.is_null())
 389     NonAliasGroup = True;
 390
 391   return NonAliasGroup;
 392 }
 393
 394 isl::ast_expr IslAst::buildRunCondition(Scop &S, const isl::ast_build &Build) {
 395   isl::ast_expr RunCondition;
 396
 397   // The conditions that need to be checked at run-time for this scop are
 398   // available as an isl_set in the runtime check context from which we can
 399   // directly derive a run-time condition.
 400   auto PosCond = Build.expr_from(S.getAssumedContext());
 401   if (S.hasTrivialInvalidContext()) {
 402     RunCondition = std::move(PosCond);
 403   } else {
 404     auto ZeroV = isl::val::zero(Build.ctx());
 405     auto NegCond = Build.expr_from(S.getInvalidContext());
 406     auto NotNegCond =
 407         isl::ast_expr::from_val(std::move(ZeroV)).eq(std::move(NegCond));
 408     RunCondition =
 409         isl::manage(isl_ast_expr_and(PosCond.release(), NotNegCond.release()));
 410   }
 411
 412   // Create the alias checks from the minimal/maximal accesses in each alias
 413   // group which consists of read only and non read only (read write) accesses.
 414   // This operation is by construction quadratic in the read-write pointers and
 415   // linear in the read only pointers in each alias group.
 416   for (const Scop::MinMaxVectorPairTy &MinMaxAccessPair : S.getAliasGroups()) {
 417     auto &MinMaxReadWrite = MinMaxAccessPair.first;
 418     auto &MinMaxReadOnly = MinMaxAccessPair.second;
 419     auto RWAccEnd = MinMaxReadWrite.end();
 420
 421     for (auto RWAccIt0 = MinMaxReadWrite.begin(); RWAccIt0 != RWAccEnd;
 422          ++RWAccIt0) {
 423       for (auto RWAccIt1 = RWAccIt0 + 1; RWAccIt1 != RWAccEnd; ++RWAccIt1)
 424         RunCondition = isl::manage(isl_ast_expr_and(
 425             RunCondition.release(),
 426             buildCondition(S, Build, RWAccIt0, RWAccIt1).release()));
 427       for (const Scop::MinMaxAccessTy &ROAccIt : MinMaxReadOnly)
 428         RunCondition = isl::manage(isl_ast_expr_and(
 429             RunCondition.release(),
 430             buildCondition(S, Build, RWAccIt0, &ROAccIt).release()));
 431     }
 432   }
 433
 434   return RunCondition;
 435 }
 436
 437 /// Simple cost analysis for a given SCoP.
 438 ///
 439 /// TODO: Improve this analysis and extract it to make it usable in other
 440 ///       places too.
 441 ///       In order to improve the cost model we could either keep track of
 442 ///       performed optimizations (e.g., tiling) or compute properties on the
 443 ///       original as well as optimized SCoP (e.g., #stride-one-accesses).
 444 static bool benefitsFromPolly(Scop &Scop, bool PerformParallelTest) {
 445   if (PollyProcessUnprofitable)
 446     return true;
 447
 448   // Check if nothing interesting happened.
 449   if (!PerformParallelTest && !Scop.isOptimized() &&
 450       Scop.getAliasGroups().empty())
 451     return false;
 452
 453   // The default assumption is that Polly improves the code.
 454   return true;
 455 }
 456
 457 /// Collect statistics for the syntax tree rooted at @p Ast.
 458 static void walkAstForStatistics(const isl::ast_node &Ast) {
 459   assert(!Ast.is_null());
 460   isl_ast_node_foreach_descendant_top_down(
 461       Ast.get(),
 462       [](__isl_keep isl_ast_node *Node, void *User) -> isl_bool {
 463         switch (isl_ast_node_get_type(Node)) {
 464         case isl_ast_node_for:
 465           NumForLoops++;
 466           if (IslAstInfo::isParallel(isl::manage_copy(Node)))
 467             NumParallel++;
 468           if (IslAstInfo::isInnermostParallel(isl::manage_copy(Node)))
 469             NumInnermostParallel++;
 470           if (IslAstInfo::isOutermostParallel(isl::manage_copy(Node)))
 471             NumOutermostParallel++;
 472           if (IslAstInfo::isReductionParallel(isl::manage_copy(Node)))
 473             NumReductionParallel++;
 474           if (IslAstInfo::isExecutedInParallel(isl::manage_copy(Node)))
 475             NumExecutedInParallel++;
 476           break;
 477
 478         case isl_ast_node_if:
 479           NumIfConditions++;
 480           break;
 481
 482         default:
 483           break;
 484         }
 485
 486         // Continue traversing subtrees.
 487         return isl_bool_true;
 488       },
 489       nullptr);
 490 }
 491
 492 IslAst::IslAst(Scop &Scop) : S(Scop), Ctx(Scop.getSharedIslCtx()) {}
 493
 494 IslAst::IslAst(IslAst &&O)
 495     : S(O.S), Ctx(O.Ctx), RunCondition(std::move(O.RunCondition)),
 496       Root(std::move(O.Root)) {}
 497
 498 void IslAst::init(const Dependences &D) {
 499   bool PerformParallelTest = PollyParallel || DetectParallel ||
 500                              PollyVectorizerChoice != VECTORIZER_NONE;
 501   auto ScheduleTree = S.getScheduleTree();
 502
 503   // Skip AST and code generation if there was no benefit achieved.
 504   if (!benefitsFromPolly(S, PerformParallelTest))
 505     return;
 506
 507   auto ScopStats = S.getStatistics();
 508   ScopsBeneficial++;
 509   BeneficialAffineLoops += ScopStats.NumAffineLoops;
 510   BeneficialBoxedLoops += ScopStats.NumBoxedLoops;
 511
 512   auto Ctx = S.getIslCtx();
 513   isl_options_set_ast_build_atomic_upper_bound(Ctx.get(), true);
 514   isl_options_set_ast_build_detect_min_max(Ctx.get(), true);
 515   isl_ast_build *Build;
 516   AstBuildUserInfo BuildInfo;
 517
 518   if (UseContext)
 519     Build = isl_ast_build_from_context(S.getContext().release());
 520   else
 521     Build = isl_ast_build_from_context(
 522         isl_set_universe(S.getParamSpace().release()));
 523
 524   Build = isl_ast_build_set_at_each_domain(Build, AtEachDomain, nullptr);
 525
 526   if (PerformParallelTest) {
 527     BuildInfo.Deps = &D;
 528     BuildInfo.InParallelFor = false;
 529     BuildInfo.InSIMD = false;
 530
 531     Build = isl_ast_build_set_before_each_for(Build, &astBuildBeforeFor,
 532                                               &BuildInfo);
 533     Build =
 534         isl_ast_build_set_after_each_for(Build, &astBuildAfterFor, &BuildInfo);
 535
 536     Build = isl_ast_build_set_before_each_mark(Build, &astBuildBeforeMark,
 537                                                &BuildInfo);
 538
 539     Build = isl_ast_build_set_after_each_mark(Build, &astBuildAfterMark,
 540                                               &BuildInfo);
 541   }
 542
 543   RunCondition = buildRunCondition(S, isl::manage_copy(Build));
 544
 545   Root = isl::manage(
 546       isl_ast_build_node_from_schedule(Build, S.getScheduleTree().release()));
 547   walkAstForStatistics(Root);
 548
 549   isl_ast_build_free(Build);
 550 }
 551
 552 IslAst IslAst::create(Scop &Scop, const Dependences &D) {
 553   IslAst Ast{Scop};
 554   Ast.init(D);
 555   return Ast;
 556 }
 557
 558 isl::ast_node IslAst::getAst() { return Root; }
 559 isl::ast_expr IslAst::getRunCondition() { return RunCondition; }
 560
 561 isl::ast_node IslAstInfo::getAst() { return Ast.getAst(); }
 562 isl::ast_expr IslAstInfo::getRunCondition() { return Ast.getRunCondition(); }
 563
 564 IslAstUserPayload *IslAstInfo::getNodePayload(const isl::ast_node &Node) {
 565   isl::id Id = Node.get_annotation();
 566   if (Id.is_null())
 567     return nullptr;
 568   IslAstUserPayload *Payload = (IslAstUserPayload *)Id.get_user();
 569   return Payload;
 570 }
 571
 572 bool IslAstInfo::isInnermost(const isl::ast_node &Node) {
 573   IslAstUserPayload *Payload = getNodePayload(Node);
 574   return Payload && Payload->IsInnermost;
 575 }
 576
 577 bool IslAstInfo::isParallel(const isl::ast_node &Node) {
 578   return IslAstInfo::isInnermostParallel(Node) ||
 579          IslAstInfo::isOutermostParallel(Node);
 580 }
 581
 582 bool IslAstInfo::isInnermostParallel(const isl::ast_node &Node) {
 583   IslAstUserPayload *Payload = getNodePayload(Node);
 584   return Payload && Payload->IsInnermostParallel;
 585 }
 586
 587 bool IslAstInfo::isOutermostParallel(const isl::ast_node &Node) {
 588   IslAstUserPayload *Payload = getNodePayload(Node);
 589   return Payload && Payload->IsOutermostParallel;
 590 }
 591
 592 bool IslAstInfo::isReductionParallel(const isl::ast_node &Node) {
 593   IslAstUserPayload *Payload = getNodePayload(Node);
 594   return Payload && Payload->IsReductionParallel;
 595 }
 596
 597 bool IslAstInfo::isExecutedInParallel(const isl::ast_node &Node) {
 598   if (!PollyParallel)
 599     return false;
 600
 601   // Do not parallelize innermost loops.
 602   //
 603   // Parallelizing innermost loops is often not profitable, especially if
 604   // they have a low number of iterations.
 605   //
 606   // TODO: Decide this based on the number of loop iterations that will be
 607   //       executed. This can possibly require run-time checks, which again
 608   //       raises the question of both run-time check overhead and code size
 609   //       costs.
 610   if (!PollyParallelForce && isInnermost(Node))
 611     return false;
 612
 613   return isOutermostParallel(Node) && !isReductionParallel(Node);
 614 }
 615
 616 isl::union_map IslAstInfo::getSchedule(const isl::ast_node &Node) {
 617   IslAstUserPayload *Payload = getNodePayload(Node);
 618   return Payload ? Payload->Build.get_schedule() : isl::union_map();
 619 }
 620
 621 isl::pw_aff
 622 IslAstInfo::getMinimalDependenceDistance(const isl::ast_node &Node) {
 623   IslAstUserPayload *Payload = getNodePayload(Node);
 624   return Payload ? Payload->MinimalDependenceDistance : isl::pw_aff();
 625 }
 626
 627 IslAstInfo::MemoryAccessSet *
 628 IslAstInfo::getBrokenReductions(const isl::ast_node &Node) {
 629   IslAstUserPayload *Payload = getNodePayload(Node);
 630   return Payload ? &Payload->BrokenReductions : nullptr;
 631 }
 632
 633 isl::ast_build IslAstInfo::getBuild(const isl::ast_node &Node) {
 634   IslAstUserPayload *Payload = getNodePayload(Node);
 635   return Payload ? Payload->Build : isl::ast_build();
 636 }
 637
 638 static std::unique_ptr<IslAstInfo> runIslAst(
 639     Scop &Scop,
 640     function_ref<const Dependences &(Dependences::AnalysisLevel)> GetDeps) {
 641   ScopsProcessed++;
 642
 643   const Dependences &D = GetDeps(Dependences::AL_Statement);
 644
 645   if (D.getSharedIslCtx() != Scop.getSharedIslCtx()) {
 646     LLVM_DEBUG(
 647         dbgs() << "Got dependence analysis for different SCoP/isl_ctx\n");
 648     return {};
 649   }
 650
 651   std::unique_ptr<IslAstInfo> Ast = std::make_unique<IslAstInfo>(Scop, D);
 652
 653   LLVM_DEBUG({
 654     if (Ast)
 655       Ast->print(dbgs());
 656   });
 657
 658   return Ast;
 659 }
 660
 661 IslAstInfo IslAstAnalysis::run(Scop &S, ScopAnalysisManager &SAM,
 662                                ScopStandardAnalysisResults &SAR) {
 663   auto GetDeps = [&](Dependences::AnalysisLevel Lvl) -> const Dependences & {
 664     return SAM.getResult<DependenceAnalysis>(S, SAR).getDependences(Lvl);
 665   };
 666
 667   return std::move(*runIslAst(S, GetDeps));
 668 }
 669
 670 static __isl_give isl_printer *cbPrintUser(__isl_take isl_printer *P,
 671                                            __isl_take isl_ast_print_options *O,
 672                                            __isl_keep isl_ast_node *Node,
 673                                            void *User) {
 674   isl::ast_node_user AstNode = isl::manage_copy(Node).as<isl::ast_node_user>();
 675   isl::ast_expr NodeExpr = AstNode.expr();
 676   isl::ast_expr CallExpr = NodeExpr.get_op_arg(0);
 677   isl::id CallExprId = CallExpr.get_id();
 678   ScopStmt *AccessStmt = (ScopStmt *)CallExprId.get_user();
 679
 680   P = isl_printer_start_line(P);
 681   P = isl_printer_print_str(P, AccessStmt->getBaseName());
 682   P = isl_printer_print_str(P, "(");
 683   P = isl_printer_end_line(P);
 684   P = isl_printer_indent(P, 2);
 685
 686   for (MemoryAccess *MemAcc : *AccessStmt) {
 687     P = isl_printer_start_line(P);
 688
 689     if (MemAcc->isRead())
 690       P = isl_printer_print_str(P, "/* read  */ &");
 691     else
 692       P = isl_printer_print_str(P, "/* write */  ");
 693
 694     isl::ast_build Build = IslAstInfo::getBuild(isl::manage_copy(Node));
 695     if (MemAcc->isAffine()) {
 696       isl_pw_multi_aff *PwmaPtr =
 697           MemAcc->applyScheduleToAccessRelation(Build.get_schedule()).release();
 698       isl::pw_multi_aff Pwma = isl::manage(PwmaPtr);
 699       isl::ast_expr AccessExpr = Build.access_from(Pwma);
 700       P = isl_printer_print_ast_expr(P, AccessExpr.get());
 701     } else {
 702       P = isl_printer_print_str(
 703           P, MemAcc->getLatestScopArrayInfo()->getName().c_str());
 704       P = isl_printer_print_str(P, "[*]");
 705     }
 706     P = isl_printer_end_line(P);
 707   }
 708
 709   P = isl_printer_indent(P, -2);
 710   P = isl_printer_start_line(P);
 711   P = isl_printer_print_str(P, ");");
 712   P = isl_printer_end_line(P);
 713
 714   isl_ast_print_options_free(O);
 715   return P;
 716 }
 717
 718 void IslAstInfo::print(raw_ostream &OS) {
 719   isl_ast_print_options *Options;
 720   isl::ast_node RootNode = Ast.getAst();
 721   Function &F = S.getFunction();
 722
 723   OS << ":: isl ast :: " << F.getName() << " :: " << S.getNameStr() << "\n";
 724
 725   if (RootNode.is_null()) {
 726     OS << ":: isl ast generation and code generation was skipped!\n\n";
 727     OS << ":: This is either because no useful optimizations could be applied "
 728           "(use -polly-process-unprofitable to enforce code generation) or "
 729           "because earlier passes such as dependence analysis timed out (use "
 730           "-polly-dependences-computeout=0 to set dependence analysis timeout "
 731           "to infinity)\n\n";
 732     return;
 733   }
 734
 735   isl::ast_expr RunCondition = Ast.getRunCondition();
 736   char *RtCStr, *AstStr;
 737
 738   Options = isl_ast_print_options_alloc(S.getIslCtx().get());
 739
 740   if (PrintAccesses)
 741     Options =
 742         isl_ast_print_options_set_print_user(Options, cbPrintUser, nullptr);
 743   Options = isl_ast_print_options_set_print_for(Options, cbPrintFor, nullptr);
 744
 745   isl_printer *P = isl_printer_to_str(S.getIslCtx().get());
 746   P = isl_printer_set_output_format(P, ISL_FORMAT_C);
 747   P = isl_printer_print_ast_expr(P, RunCondition.get());
 748   RtCStr = isl_printer_get_str(P);
 749   P = isl_printer_flush(P);
 750   P = isl_printer_indent(P, 4);
 751   P = isl_ast_node_print(RootNode.get(), P, Options);
 752   AstStr = isl_printer_get_str(P);
 753
 754   LLVM_DEBUG({
 755     dbgs() << S.getContextStr() << "\n";
 756     dbgs() << stringFromIslObj(S.getScheduleTree(), "null");
 757   });
 758   OS << "\nif (" << RtCStr << ")\n\n";
 759   OS << AstStr << "\n";
 760   OS << "else\n";
 761   OS << "    {  /* original code */ }\n\n";
 762
 763   free(RtCStr);
 764   free(AstStr);
 765
 766   isl_printer_free(P);
 767 }
 768
 769 AnalysisKey IslAstAnalysis::Key;
 770 PreservedAnalyses IslAstPrinterPass::run(Scop &S, ScopAnalysisManager &SAM,
 771                                          ScopStandardAnalysisResults &SAR,
 772                                          SPMUpdater &U) {
 773   auto &Ast = SAM.getResult<IslAstAnalysis>(S, SAR);
 774   Ast.print(OS);
 775   return PreservedAnalyses::all();
 776 }
 777
 778 void IslAstInfoWrapperPass::releaseMemory() { Ast.reset(); }
 779
 780 bool IslAstInfoWrapperPass::runOnScop(Scop &Scop) {
 781   auto GetDeps = [this](Dependences::AnalysisLevel Lvl) -> const Dependences & {
 782     return getAnalysis<DependenceInfo>().getDependences(Lvl);
 783   };
 784
 785   Ast = runIslAst(Scop, GetDeps);
 786
 787   return false;
 788 }
 789
 790 void IslAstInfoWrapperPass::getAnalysisUsage(AnalysisUsage &AU) const {
 791   // Get the Common analysis usage of ScopPasses.
 792   ScopPass::getAnalysisUsage(AU);
 793   AU.addRequiredTransitive<ScopInfoRegionPass>();
 794   AU.addRequired<DependenceInfo>();
 795
 796   AU.addPreserved<DependenceInfo>();
 797 }
 798
 799 void IslAstInfoWrapperPass::printScop(raw_ostream &OS, Scop &S) const {
 800   OS << "Printing analysis 'Polly - Generate an AST of the SCoP (isl)'"
 801      << S.getName() << "' in function '" << S.getFunction().getName() << "':\n";
 802   if (Ast)
 803     Ast->print(OS);
 804 }
 805
 806 char IslAstInfoWrapperPass::ID = 0;
 807
 808 Pass *polly::createIslAstInfoWrapperPassPass() {
 809   return new IslAstInfoWrapperPass();
 810 }
 811
 812 INITIALIZE_PASS_BEGIN(IslAstInfoWrapperPass, "polly-ast",
 813                       "Polly - Generate an AST of the SCoP (isl)", false,
 814                       false);
 815 INITIALIZE_PASS_DEPENDENCY(ScopInfoRegionPass);
 816 INITIALIZE_PASS_DEPENDENCY(DependenceInfo);
 817 INITIALIZE_PASS_END(IslAstInfoWrapperPass, "polly-ast",
 818                     "Polly - Generate an AST from the SCoP (isl)", false, false)
 819
 820 //===----------------------------------------------------------------------===//
 821
 822 namespace {
 823 /// Print result from IslAstInfoWrapperPass.
 824 class IslAstInfoPrinterLegacyPass final : public ScopPass {
 825 public:
 826   static char ID;
 827
 828   IslAstInfoPrinterLegacyPass() : IslAstInfoPrinterLegacyPass(outs()) {}
 829   explicit IslAstInfoPrinterLegacyPass(llvm::raw_ostream &OS)
 830       : ScopPass(ID), OS(OS) {}
 831
 832   bool runOnScop(Scop &S) override {
 833     IslAstInfoWrapperPass &P = getAnalysis<IslAstInfoWrapperPass>();
 834
 835     OS << "Printing analysis '" << P.getPassName() << "' for region: '"
 836        << S.getRegion().getNameStr() << "' in function '"
 837        << S.getFunction().getName() << "':\n";
 838     P.printScop(OS, S);
 839
 840     return false;
 841   }
 842
 843   void getAnalysisUsage(AnalysisUsage &AU) const override {
 844     ScopPass::getAnalysisUsage(AU);
 845     AU.addRequired<IslAstInfoWrapperPass>();
 846     AU.setPreservesAll();
 847   }
 848
 849 private:
 850   llvm::raw_ostream &OS;
 851 };
 852
 853 char IslAstInfoPrinterLegacyPass::ID = 0;
 854 } // namespace
 855
 856 Pass *polly::createIslAstInfoPrinterLegacyPass(raw_ostream &OS) {
 857   return new IslAstInfoPrinterLegacyPass(OS);
 858 }
 859
 860 INITIALIZE_PASS_BEGIN(IslAstInfoPrinterLegacyPass, "polly-print-ast",
 861                       "Polly - Print the AST from a SCoP (isl)", false, false);
 862 INITIALIZE_PASS_DEPENDENCY(IslAstInfoWrapperPass);
 863 INITIALIZE_PASS_END(IslAstInfoPrinterLegacyPass, "polly-print-ast",
 864                     "Polly - Print the AST from a SCoP (isl)", false, false)