1 //===- IslAst.cpp - isl code generator interface --------------------------===//
2 //
3 //                     The LLVM Compiler Infrastructure
4 //
5 // This file is distributed under the University of Illinois Open Source
6 // License. See LICENSE.TXT for details.
7 //
8 //===----------------------------------------------------------------------===//
9 //
10 // The isl code generator interface takes a Scop and generates a isl_ast. This
11 // ist_ast can either be returned directly or it can be pretty printed to
12 // stdout.
13 //
14 // A typical isl_ast output looks like this:
15 //
16 // for (c2 = max(0, ceild(n + m, 2); c2 <= min(511, floord(5 * n, 3)); c2++) {
17 //   bb2(c2);
18 // }
19 //
20 //===----------------------------------------------------------------------===//
21 
22 #include "polly/CodeGen/CodeGeneration.h"
23 #include "polly/CodeGen/IslAst.h"
24 #include "polly/Dependences.h"
25 #include "polly/LinkAllPasses.h"
26 #include "polly/Options.h"
27 #include "polly/ScopInfo.h"
28 #include "llvm/Support/Debug.h"
29 
30 #include "isl/union_map.h"
31 #include "isl/list.h"
32 #include "isl/ast_build.h"
33 #include "isl/set.h"
34 #include "isl/map.h"
35 #include "isl/aff.h"
36 
37 #define DEBUG_TYPE "polly-ast"
38 
39 using namespace llvm;
40 using namespace polly;
41 
42 using IslAstUserPayload = IslAstInfo::IslAstUserPayload;
43 
44 static cl::opt<bool> UseContext("polly-ast-use-context",
45                                 cl::desc("Use context"), cl::Hidden,
46                                 cl::init(false), cl::ZeroOrMore,
47                                 cl::cat(PollyCategory));
48 
49 static cl::opt<bool> DetectParallel("polly-ast-detect-parallel",
50                                     cl::desc("Detect parallelism"), cl::Hidden,
51                                     cl::init(false), cl::ZeroOrMore,
52                                     cl::cat(PollyCategory));
53 
54 namespace polly {
55 class IslAst {
56 public:
57   IslAst(Scop *Scop, Dependences &D);
58 
59   ~IslAst();
60 
61   /// Print a source code representation of the program.
62   void pprint(llvm::raw_ostream &OS);
63 
64   __isl_give isl_ast_node *getAst();
65 
66   /// @brief Get the run-time conditions for the Scop.
67   __isl_give isl_ast_expr *getRunCondition();
68 
69 private:
70   Scop *S;
71   isl_ast_node *Root;
72   isl_ast_expr *RunCondition;
73 
74   void buildRunCondition(__isl_keep isl_ast_build *Build);
75 };
76 } // End namespace polly.
77 
78 /// @brief Free an IslAstUserPayload object pointed to by @p Ptr
79 static void freeIslAstUserPayload(void *Ptr) {
80   delete ((IslAstInfo::IslAstUserPayload *)Ptr);
81 }
82 
83 IslAstInfo::IslAstUserPayload::~IslAstUserPayload() {
84   isl_ast_build_free(Build);
85   isl_pw_aff_free(MinimalDependenceDistance);
86 }
87 
88 /// @brief Temporary information used when building the ast.
89 struct AstBuildUserInfo {
90   /// @brief Construct and initialize the helper struct for AST creation.
91   AstBuildUserInfo()
92       : Deps(nullptr), InParallelFor(false), LastForNodeId(nullptr) {}
93 
94   /// @brief The dependence information used for the parallelism check.
95   Dependences *Deps;
96 
97   /// @brief Flag to indicate that we are inside a parallel for node.
98   bool InParallelFor;
99 
100   /// @brief The last iterator id created for the current SCoP.
101   isl_id *LastForNodeId;
102 };
103 
104 /// @brief Print a string @p str in a single line using @p Printer.
105 static isl_printer *printLine(__isl_take isl_printer *Printer,
106                               const std::string &str,
107                               __isl_keep isl_pw_aff *PWA = nullptr) {
108   Printer = isl_printer_start_line(Printer);
109   Printer = isl_printer_print_str(Printer, str.c_str());
110   if (PWA)
111     Printer = isl_printer_print_pw_aff(Printer, PWA);
112   return isl_printer_end_line(Printer);
113 }
114 
115 /// @brief Return all broken reductions as a string of clauses (OpenMP style).
116 static const std::string getBrokenReductionsStr(__isl_keep isl_ast_node *Node) {
117   IslAstInfo::MemoryAccessSet *BrokenReductions;
118   std::string str;
119 
120   BrokenReductions = IslAstInfo::getBrokenReductions(Node);
121   if (!BrokenReductions || BrokenReductions->empty())
122     return "";
123 
124   // Map each type of reduction to a comma separated list of the base addresses.
125   std::map<MemoryAccess::ReductionType, std::string> Clauses;
126   for (MemoryAccess *MA : *BrokenReductions)
127     if (MA->isWrite())
128       Clauses[MA->getReductionType()] +=
129           ", " + MA->getBaseAddr()->getName().str();
130 
131   // Now print the reductions sorted by type. Each type will cause a clause
132   // like:  reduction (+ : sum0, sum1, sum2)
133   for (const auto &ReductionClause : Clauses) {
134     str += " reduction (";
135     str += MemoryAccess::getReductionOperatorStr(ReductionClause.first);
136     // Remove the first two symbols (", ") to make the output look pretty.
137     str += " : " + ReductionClause.second.substr(2) + ")";
138   }
139 
140   return str;
141 }
142 
143 /// @brief Callback executed for each for node in the ast in order to print it.
144 static isl_printer *cbPrintFor(__isl_take isl_printer *Printer,
145                                __isl_take isl_ast_print_options *Options,
146                                __isl_keep isl_ast_node *Node, void *) {
147 
148   isl_pw_aff *DD = IslAstInfo::getMinimalDependenceDistance(Node);
149   const std::string BrokenReductionsStr = getBrokenReductionsStr(Node);
150   const std::string DepDisPragmaStr = "#pragma minimal dependence distance: ";
151   const std::string SimdPragmaStr = "#pragma simd";
152   const std::string OmpPragmaStr = "#pragma omp parallel for";
153 
154   if (DD)
155     Printer = printLine(Printer, DepDisPragmaStr, DD);
156 
157   if (IslAstInfo::isInnermostParallel(Node))
158     Printer = printLine(Printer, SimdPragmaStr + BrokenReductionsStr);
159 
160   if (IslAstInfo::isOutermostParallel(Node))
161     Printer = printLine(Printer, OmpPragmaStr + BrokenReductionsStr);
162 
163   isl_pw_aff_free(DD);
164   return isl_ast_node_for_print(Node, Printer, Options);
165 }
166 
167 /// @brief Check if the current scheduling dimension is parallel
168 ///
169 /// In case the dimension is parallel we also check if any reduction
170 /// dependences is broken when we exploit this parallelism. If so,
171 /// @p IsReductionParallel will be set to true. The reduction dependences we use
172 /// to check are actually the union of the transitive closure of the initial
173 /// reduction dependences together with their reveresal. Even though these
174 /// dependences connect all iterations with each other (thus they are cyclic)
175 /// we can perform the parallelism check as we are only interested in a zero
176 /// (or non-zero) dependence distance on the dimension in question.
177 static bool astScheduleDimIsParallel(__isl_keep isl_ast_build *Build,
178                                      Dependences *D,
179                                      IslAstUserPayload *NodeInfo) {
180   if (!D->hasValidDependences())
181     return false;
182 
183   isl_union_map *Schedule = isl_ast_build_get_schedule(Build);
184   isl_union_map *Deps = D->getDependences(
185       Dependences::TYPE_RAW | Dependences::TYPE_WAW | Dependences::TYPE_WAR);
186 
187   if (!D->isParallel(Schedule, Deps, &NodeInfo->MinimalDependenceDistance) &&
188       !isl_union_map_free(Schedule))
189     return false;
190 
191   isl_union_map *RedDeps = D->getDependences(Dependences::TYPE_TC_RED);
192   if (!D->isParallel(Schedule, RedDeps))
193     NodeInfo->IsReductionParallel = true;
194 
195   if (!NodeInfo->IsReductionParallel && !isl_union_map_free(Schedule))
196     return true;
197 
198   // Annotate reduction parallel nodes with the memory accesses which caused the
199   // reduction dependences parallel execution of the node conflicts with.
200   for (const auto &MaRedPair : D->getReductionDependences()) {
201     if (!MaRedPair.second)
202       continue;
203     RedDeps = isl_union_map_from_map(isl_map_copy(MaRedPair.second));
204     if (!D->isParallel(Schedule, RedDeps))
205       NodeInfo->BrokenReductions.insert(MaRedPair.first);
206   }
207 
208   isl_union_map_free(Schedule);
209   return true;
210 }
211 
212 // This method is executed before the construction of a for node. It creates
213 // an isl_id that is used to annotate the subsequently generated ast for nodes.
214 //
215 // In this function we also run the following analyses:
216 //
217 // - Detection of openmp parallel loops
218 //
219 static __isl_give isl_id *astBuildBeforeFor(__isl_keep isl_ast_build *Build,
220                                             void *User) {
221   AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
222   IslAstUserPayload *Payload = new IslAstUserPayload();
223   isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", Payload);
224   Id = isl_id_set_free_user(Id, freeIslAstUserPayload);
225   BuildInfo->LastForNodeId = Id;
226 
227   // Test for parallelism only if we are not already inside a parallel loop
228   if (!BuildInfo->InParallelFor)
229     BuildInfo->InParallelFor = Payload->IsOutermostParallel =
230         astScheduleDimIsParallel(Build, BuildInfo->Deps, Payload);
231 
232   return Id;
233 }
234 
235 // This method is executed after the construction of a for node.
236 //
237 // It performs the following actions:
238 //
239 // - Reset the 'InParallelFor' flag, as soon as we leave a for node,
240 //   that is marked as openmp parallel.
241 //
242 static __isl_give isl_ast_node *
243 astBuildAfterFor(__isl_take isl_ast_node *Node, __isl_keep isl_ast_build *Build,
244                  void *User) {
245   isl_id *Id = isl_ast_node_get_annotation(Node);
246   assert(Id && "Post order visit assumes annotated for nodes");
247   IslAstUserPayload *Payload = (IslAstUserPayload *)isl_id_get_user(Id);
248   assert(Payload && "Post order visit assumes annotated for nodes");
249 
250   AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
251   assert(!Payload->Build && "Build environment already set");
252   Payload->Build = isl_ast_build_copy(Build);
253   Payload->IsInnermost = (Id == BuildInfo->LastForNodeId);
254 
255   // Innermost loops that are surrounded by parallel loops have not yet been
256   // tested for parallelism. Test them here to ensure we check all innermost
257   // loops for parallelism.
258   if (Payload->IsInnermost && BuildInfo->InParallelFor) {
259     if (Payload->IsOutermostParallel)
260       Payload->IsInnermostParallel = true;
261     else
262       Payload->IsInnermostParallel =
263           astScheduleDimIsParallel(Build, BuildInfo->Deps, Payload);
264   }
265   if (Payload->IsOutermostParallel)
266     BuildInfo->InParallelFor = false;
267 
268   isl_id_free(Id);
269   return Node;
270 }
271 
272 static __isl_give isl_ast_node *AtEachDomain(__isl_take isl_ast_node *Node,
273                                              __isl_keep isl_ast_build *Build,
274                                              void *User) {
275   assert(!isl_ast_node_get_annotation(Node) && "Node already annotated");
276 
277   IslAstUserPayload *Payload = new IslAstUserPayload();
278   isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", Payload);
279   Id = isl_id_set_free_user(Id, freeIslAstUserPayload);
280 
281   Payload->Build = isl_ast_build_copy(Build);
282 
283   return isl_ast_node_set_annotation(Node, Id);
284 }
285 
286 void IslAst::buildRunCondition(__isl_keep isl_ast_build *Build) {
287   // The conditions that need to be checked at run-time for this scop are
288   // available as an isl_set in the AssumedContext. We generate code for this
289   // check as follows. First, we generate an isl_pw_aff that is 1, if a certain
290   // combination of parameter values fulfills the conditions in the assumed
291   // context, and that is 0 otherwise. We then translate this isl_pw_aff into
292   // an isl_ast_expr. At run-time this expression can be evaluated and the
293   // optimized scop can be executed conditionally according to the result of the
294   // run-time check.
295 
296   isl_aff *Zero =
297       isl_aff_zero_on_domain(isl_local_space_from_space(S->getParamSpace()));
298   isl_aff *One =
299       isl_aff_zero_on_domain(isl_local_space_from_space(S->getParamSpace()));
300 
301   One = isl_aff_add_constant_si(One, 1);
302 
303   isl_pw_aff *PwZero = isl_pw_aff_from_aff(Zero);
304   isl_pw_aff *PwOne = isl_pw_aff_from_aff(One);
305 
306   PwOne = isl_pw_aff_intersect_domain(PwOne, S->getAssumedContext());
307   PwZero = isl_pw_aff_intersect_domain(
308       PwZero, isl_set_complement(S->getAssumedContext()));
309 
310   isl_pw_aff *Cond = isl_pw_aff_union_max(PwOne, PwZero);
311 
312   RunCondition = isl_ast_build_expr_from_pw_aff(Build, Cond);
313 
314   // Create the alias checks from the minimal/maximal accesses in each alias
315   // group. This operation is by construction quadratic in the number of
316   // elements in each alias group.
317   isl_ast_expr *NonAliasGroup, *MinExpr, *MaxExpr;
318   for (const Scop::MinMaxVectorTy *MinMaxAccesses : S->getAliasGroups()) {
319     auto AccEnd = MinMaxAccesses->end();
320     for (auto AccIt0 = MinMaxAccesses->begin(); AccIt0 != AccEnd; ++AccIt0) {
321       for (auto AccIt1 = AccIt0 + 1; AccIt1 != AccEnd; ++AccIt1) {
322         MinExpr =
323             isl_ast_expr_address_of(isl_ast_build_access_from_pw_multi_aff(
324                 Build, isl_pw_multi_aff_copy(AccIt0->first)));
325         MaxExpr =
326             isl_ast_expr_address_of(isl_ast_build_access_from_pw_multi_aff(
327                 Build, isl_pw_multi_aff_copy(AccIt1->second)));
328         NonAliasGroup = isl_ast_expr_le(MaxExpr, MinExpr);
329         MinExpr =
330             isl_ast_expr_address_of(isl_ast_build_access_from_pw_multi_aff(
331                 Build, isl_pw_multi_aff_copy(AccIt1->first)));
332         MaxExpr =
333             isl_ast_expr_address_of(isl_ast_build_access_from_pw_multi_aff(
334                 Build, isl_pw_multi_aff_copy(AccIt0->second)));
335         NonAliasGroup =
336             isl_ast_expr_or(NonAliasGroup, isl_ast_expr_le(MaxExpr, MinExpr));
337         RunCondition = isl_ast_expr_and(RunCondition, NonAliasGroup);
338       }
339     }
340   }
341 }
342 
343 IslAst::IslAst(Scop *Scop, Dependences &D) : S(Scop) {
344   isl_ctx *Ctx = S->getIslCtx();
345   isl_options_set_ast_build_atomic_upper_bound(Ctx, true);
346   isl_ast_build *Build;
347   AstBuildUserInfo BuildInfo;
348 
349   if (UseContext)
350     Build = isl_ast_build_from_context(S->getContext());
351   else
352     Build = isl_ast_build_from_context(isl_set_universe(S->getParamSpace()));
353 
354   Build = isl_ast_build_set_at_each_domain(Build, AtEachDomain, nullptr);
355 
356   isl_union_map *Schedule =
357       isl_union_map_intersect_domain(S->getSchedule(), S->getDomains());
358 
359   if (DetectParallel || PollyVectorizerChoice != VECTORIZER_NONE) {
360     BuildInfo.Deps = &D;
361     BuildInfo.InParallelFor = 0;
362 
363     Build = isl_ast_build_set_before_each_for(Build, &astBuildBeforeFor,
364                                               &BuildInfo);
365     Build =
366         isl_ast_build_set_after_each_for(Build, &astBuildAfterFor, &BuildInfo);
367   }
368 
369   buildRunCondition(Build);
370 
371   Root = isl_ast_build_ast_from_schedule(Build, Schedule);
372 
373   isl_ast_build_free(Build);
374 }
375 
376 IslAst::~IslAst() {
377   isl_ast_node_free(Root);
378   isl_ast_expr_free(RunCondition);
379 }
380 
381 __isl_give isl_ast_node *IslAst::getAst() { return isl_ast_node_copy(Root); }
382 __isl_give isl_ast_expr *IslAst::getRunCondition() {
383   return isl_ast_expr_copy(RunCondition);
384 }
385 
386 void IslAstInfo::releaseMemory() {
387   if (Ast) {
388     delete Ast;
389     Ast = 0;
390   }
391 }
392 
393 bool IslAstInfo::runOnScop(Scop &Scop) {
394   if (Ast)
395     delete Ast;
396 
397   S = &Scop;
398 
399   Dependences &D = getAnalysis<Dependences>();
400 
401   Ast = new IslAst(&Scop, D);
402 
403   DEBUG(printScop(dbgs()));
404   return false;
405 }
406 
407 __isl_give isl_ast_node *IslAstInfo::getAst() const { return Ast->getAst(); }
408 __isl_give isl_ast_expr *IslAstInfo::getRunCondition() const {
409   return Ast->getRunCondition();
410 }
411 
412 IslAstUserPayload *IslAstInfo::getNodePayload(__isl_keep isl_ast_node *Node) {
413   isl_id *Id = isl_ast_node_get_annotation(Node);
414   if (!Id)
415     return nullptr;
416   IslAstUserPayload *Payload = (IslAstUserPayload *)isl_id_get_user(Id);
417   isl_id_free(Id);
418   return Payload;
419 }
420 
421 bool IslAstInfo::isInnermost(__isl_keep isl_ast_node *Node) {
422   IslAstUserPayload *Payload = getNodePayload(Node);
423   return Payload && Payload->IsInnermost;
424 }
425 
426 bool IslAstInfo::isParallel(__isl_keep isl_ast_node *Node) {
427   return IslAstInfo::isInnermostParallel(Node) ||
428          IslAstInfo::isOutermostParallel(Node);
429 }
430 
431 bool IslAstInfo::isInnermostParallel(__isl_keep isl_ast_node *Node) {
432   IslAstUserPayload *Payload = getNodePayload(Node);
433   return Payload && Payload->IsInnermostParallel;
434 }
435 
436 bool IslAstInfo::isOutermostParallel(__isl_keep isl_ast_node *Node) {
437   IslAstUserPayload *Payload = getNodePayload(Node);
438   return Payload && Payload->IsOutermostParallel;
439 }
440 
441 bool IslAstInfo::isReductionParallel(__isl_keep isl_ast_node *Node) {
442   IslAstUserPayload *Payload = getNodePayload(Node);
443   return Payload && Payload->IsReductionParallel;
444 }
445 
446 isl_union_map *IslAstInfo::getSchedule(__isl_keep isl_ast_node *Node) {
447   IslAstUserPayload *Payload = getNodePayload(Node);
448   return Payload ? isl_ast_build_get_schedule(Payload->Build) : nullptr;
449 }
450 
451 isl_pw_aff *
452 IslAstInfo::getMinimalDependenceDistance(__isl_keep isl_ast_node *Node) {
453   IslAstUserPayload *Payload = getNodePayload(Node);
454   return Payload ? isl_pw_aff_copy(Payload->MinimalDependenceDistance)
455                  : nullptr;
456 }
457 
458 IslAstInfo::MemoryAccessSet *
459 IslAstInfo::getBrokenReductions(__isl_keep isl_ast_node *Node) {
460   IslAstUserPayload *Payload = getNodePayload(Node);
461   return Payload ? &Payload->BrokenReductions : nullptr;
462 }
463 
464 isl_ast_build *IslAstInfo::getBuild(__isl_keep isl_ast_node *Node) {
465   IslAstUserPayload *Payload = getNodePayload(Node);
466   return Payload ? Payload->Build : nullptr;
467 }
468 
469 void IslAstInfo::printScop(raw_ostream &OS) const {
470   isl_ast_print_options *Options;
471   isl_ast_node *RootNode = getAst();
472   isl_ast_expr *RunCondition = getRunCondition();
473   char *RtCStr, *AstStr;
474 
475   Scop &S = getCurScop();
476   Options = isl_ast_print_options_alloc(S.getIslCtx());
477   Options = isl_ast_print_options_set_print_for(Options, cbPrintFor, nullptr);
478 
479   isl_printer *P = isl_printer_to_str(S.getIslCtx());
480   P = isl_printer_print_ast_expr(P, RunCondition);
481   RtCStr = isl_printer_get_str(P);
482   P = isl_printer_flush(P);
483   P = isl_printer_indent(P, 4);
484   P = isl_printer_set_output_format(P, ISL_FORMAT_C);
485   P = isl_ast_node_print(RootNode, P, Options);
486   AstStr = isl_printer_get_str(P);
487 
488   Function *F = S.getRegion().getEntry()->getParent();
489   isl_union_map *Schedule =
490       isl_union_map_intersect_domain(S.getSchedule(), S.getDomains());
491 
492   OS << ":: isl ast :: " << F->getName() << " :: " << S.getRegion().getNameStr()
493      << "\n";
494   DEBUG(dbgs() << S.getContextStr() << "\n"; isl_union_map_dump(Schedule));
495   OS << "\nif (" << RtCStr << ")\n\n";
496   OS << AstStr << "\n";
497   OS << "else\n";
498   OS << "    {  /* original code */ }\n\n";
499 
500   isl_ast_expr_free(RunCondition);
501   isl_union_map_free(Schedule);
502   isl_ast_node_free(RootNode);
503   isl_printer_free(P);
504 }
505 
506 void IslAstInfo::getAnalysisUsage(AnalysisUsage &AU) const {
507   // Get the Common analysis usage of ScopPasses.
508   ScopPass::getAnalysisUsage(AU);
509   AU.addRequired<ScopInfo>();
510   AU.addRequired<Dependences>();
511 }
512 
513 char IslAstInfo::ID = 0;
514 
515 Pass *polly::createIslAstInfoPass() { return new IslAstInfo(); }
516 
517 INITIALIZE_PASS_BEGIN(IslAstInfo, "polly-ast",
518                       "Polly - Generate an AST of the SCoP (isl)", false,
519                       false);
520 INITIALIZE_PASS_DEPENDENCY(ScopInfo);
521 INITIALIZE_PASS_DEPENDENCY(Dependences);
522 INITIALIZE_PASS_END(IslAstInfo, "polly-ast",
523                     "Polly - Generate an AST from the SCoP (isl)", false, false)
524