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 using namespace llvm;
38 using namespace polly;
39 
40 #define DEBUG_TYPE "polly-ast"
41 
42 static cl::opt<bool> UseContext("polly-ast-use-context",
43                                 cl::desc("Use context"), cl::Hidden,
44                                 cl::init(false), cl::ZeroOrMore,
45                                 cl::cat(PollyCategory));
46 
47 static cl::opt<bool> DetectParallel("polly-ast-detect-parallel",
48                                     cl::desc("Detect parallelism"), cl::Hidden,
49                                     cl::init(false), cl::ZeroOrMore,
50                                     cl::cat(PollyCategory));
51 
52 namespace polly {
53 class IslAst {
54 public:
55   IslAst(Scop *Scop, Dependences &D);
56 
57   ~IslAst();
58 
59   /// Print a source code representation of the program.
60   void pprint(llvm::raw_ostream &OS);
61 
62   __isl_give isl_ast_node *getAst();
63 
64   /// @brief Get the run-time conditions for the Scop.
65   __isl_give isl_ast_expr *getRunCondition();
66 
67 private:
68   Scop *S;
69   isl_ast_node *Root;
70   isl_ast_expr *RunCondition;
71 
72   void buildRunCondition(__isl_keep isl_ast_build *Context);
73 };
74 } // End namespace polly.
75 
76 // Temporary information used when building the ast.
77 struct AstBuildUserInfo {
78   // The dependence information.
79   Dependences *Deps;
80 
81   // We are inside a parallel for node.
82   int InParallelFor;
83 };
84 
85 // Print a loop annotated with OpenMP or vector pragmas.
86 static __isl_give isl_printer *
87 printParallelFor(__isl_keep isl_ast_node *Node, __isl_take isl_printer *Printer,
88                  __isl_take isl_ast_print_options *PrintOptions,
89                  IslAstUserPayload *Info) {
90   if (Info) {
91     if (Info->IsInnermostParallel) {
92       Printer = isl_printer_start_line(Printer);
93       Printer = isl_printer_print_str(Printer, "#pragma simd");
94       if (Info->IsReductionParallel)
95         Printer = isl_printer_print_str(Printer, " reduction");
96       Printer = isl_printer_end_line(Printer);
97     }
98     if (Info->IsOutermostParallel) {
99       Printer = isl_printer_start_line(Printer);
100       Printer = isl_printer_print_str(Printer, "#pragma omp parallel for");
101       if (Info->IsReductionParallel)
102         Printer = isl_printer_print_str(Printer, " reduction");
103       Printer = isl_printer_end_line(Printer);
104     }
105   }
106   return isl_ast_node_for_print(Node, Printer, PrintOptions);
107 }
108 
109 // Print an isl_ast_for.
110 static __isl_give isl_printer *
111 printFor(__isl_take isl_printer *Printer,
112          __isl_take isl_ast_print_options *PrintOptions,
113          __isl_keep isl_ast_node *Node, void *User) {
114   isl_id *Id = isl_ast_node_get_annotation(Node);
115   if (!Id)
116     return isl_ast_node_for_print(Node, Printer, PrintOptions);
117 
118   struct IslAstUserPayload *Info =
119       (struct IslAstUserPayload *)isl_id_get_user(Id);
120   Printer = printParallelFor(Node, Printer, PrintOptions, Info);
121   isl_id_free(Id);
122   return Printer;
123 }
124 
125 // Allocate an AstNodeInfo structure and initialize it with default values.
126 static struct IslAstUserPayload *allocateIslAstUser() {
127   struct IslAstUserPayload *NodeInfo;
128   NodeInfo =
129       (struct IslAstUserPayload *)malloc(sizeof(struct IslAstUserPayload));
130   NodeInfo->Context = 0;
131   NodeInfo->IsOutermostParallel = 0;
132   NodeInfo->IsInnermostParallel = 0;
133   NodeInfo->IsReductionParallel = false;
134   return NodeInfo;
135 }
136 
137 // Free the AstNodeInfo structure.
138 static void freeIslAstUser(void *Ptr) {
139   struct IslAstUserPayload *UserStruct = (struct IslAstUserPayload *)Ptr;
140   isl_ast_build_free(UserStruct->Context);
141   free(UserStruct);
142 }
143 
144 // Check if the current scheduling dimension is parallel.
145 //
146 // We check for parallelism by verifying that the loop does not carry any
147 // dependences.
148 //
149 // Parallelism test: if the distance is zero in all outer dimensions, then it
150 // has to be zero in the current dimension as well.
151 //
152 // Implementation: first, translate dependences into time space, then force
153 // outer dimensions to be equal. If the distance is zero in the current
154 // dimension, then the loop is parallel. The distance is zero in the current
155 // dimension if it is a subset of a map with equal values for the current
156 // dimension.
157 static bool astScheduleDimIsParallel(__isl_keep isl_ast_build *Build,
158                                      __isl_take isl_union_map *Deps) {
159   isl_union_map *Schedule;
160   isl_map *ScheduleDeps, *Test;
161   isl_space *ScheduleSpace;
162   unsigned Dimension, IsParallel;
163 
164   Schedule = isl_ast_build_get_schedule(Build);
165   ScheduleSpace = isl_ast_build_get_schedule_space(Build);
166 
167   Dimension = isl_space_dim(ScheduleSpace, isl_dim_out) - 1;
168 
169   Deps = isl_union_map_apply_range(Deps, isl_union_map_copy(Schedule));
170   Deps = isl_union_map_apply_domain(Deps, Schedule);
171 
172   if (isl_union_map_is_empty(Deps)) {
173     isl_union_map_free(Deps);
174     isl_space_free(ScheduleSpace);
175     return true;
176   }
177 
178   ScheduleDeps = isl_map_from_union_map(Deps);
179 
180   for (unsigned i = 0; i < Dimension; i++)
181     ScheduleDeps = isl_map_equate(ScheduleDeps, isl_dim_out, i, isl_dim_in, i);
182 
183   Test = isl_map_universe(isl_map_get_space(ScheduleDeps));
184   Test = isl_map_equate(Test, isl_dim_out, Dimension, isl_dim_in, Dimension);
185   IsParallel = isl_map_is_subset(ScheduleDeps, Test);
186 
187   isl_space_free(ScheduleSpace);
188   isl_map_free(Test);
189   isl_map_free(ScheduleDeps);
190 
191   return IsParallel;
192 }
193 
194 /// @brief Check if the current scheduling dimension is parallel
195 ///
196 /// In case the dimension is parallel we also check if any reduction
197 /// dependences is broken when we exploit this parallelism. If so,
198 /// @p IsReductionParallel will be set to true. The reduction dependences we use
199 /// to check are actually the union of the transitive closure of the initial
200 /// reduction dependences together with their reveresal. Even though these
201 /// dependences connect all iterations with each other (thus they are cyclic)
202 /// we can perform the parallelism check as we are only interested in a zero
203 /// (or non-zero) dependence distance on the dimension in question.
204 static bool astScheduleDimIsParallel(__isl_keep isl_ast_build *Build,
205                                      Dependences *D,
206                                      bool &IsReductionParallel) {
207   if (!D->hasValidDependences())
208     return false;
209 
210   isl_union_map *Deps = D->getDependences(
211       Dependences::TYPE_RAW | Dependences::TYPE_WAW | Dependences::TYPE_WAR);
212   if (!astScheduleDimIsParallel(Build, Deps))
213     return false;
214 
215   isl_union_map *RedDeps = D->getDependences(Dependences::TYPE_TC_RED);
216   if (!astScheduleDimIsParallel(Build, RedDeps))
217     IsReductionParallel = true;
218 
219   return true;
220 }
221 
222 // Mark a for node openmp parallel, if it is the outermost parallel for node.
223 static void markOpenmpParallel(__isl_keep isl_ast_build *Build,
224                                struct AstBuildUserInfo *BuildInfo,
225                                struct IslAstUserPayload *NodeInfo) {
226   if (BuildInfo->InParallelFor)
227     return;
228 
229   if (astScheduleDimIsParallel(Build, BuildInfo->Deps,
230                                NodeInfo->IsReductionParallel)) {
231     BuildInfo->InParallelFor = 1;
232     NodeInfo->IsOutermostParallel = 1;
233   }
234 }
235 
236 // This method is executed before the construction of a for node. It creates
237 // an isl_id that is used to annotate the subsequently generated ast for nodes.
238 //
239 // In this function we also run the following analyses:
240 //
241 // - Detection of openmp parallel loops
242 //
243 static __isl_give isl_id *astBuildBeforeFor(__isl_keep isl_ast_build *Build,
244                                             void *User) {
245   struct AstBuildUserInfo *BuildInfo = (struct AstBuildUserInfo *)User;
246   struct IslAstUserPayload *NodeInfo = allocateIslAstUser();
247   isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", NodeInfo);
248   Id = isl_id_set_free_user(Id, freeIslAstUser);
249 
250   markOpenmpParallel(Build, BuildInfo, NodeInfo);
251 
252   return Id;
253 }
254 
255 // Returns 0 when Node contains loops, otherwise returns -1. This search
256 // function uses ISL's way to iterate over lists of isl_ast_nodes with
257 // isl_ast_node_list_foreach. Please use the single argument wrapper function
258 // that returns a bool instead of using this function directly.
259 static int containsLoops(__isl_take isl_ast_node *Node, void *User) {
260   if (!Node)
261     return -1;
262 
263   switch (isl_ast_node_get_type(Node)) {
264   case isl_ast_node_for:
265     isl_ast_node_free(Node);
266     return 0;
267   case isl_ast_node_block: {
268     isl_ast_node_list *List = isl_ast_node_block_get_children(Node);
269     int Res = isl_ast_node_list_foreach(List, &containsLoops, nullptr);
270     isl_ast_node_list_free(List);
271     isl_ast_node_free(Node);
272     return Res;
273   }
274   case isl_ast_node_if: {
275     int Res = -1;
276     if (0 == containsLoops(isl_ast_node_if_get_then(Node), nullptr) ||
277         (isl_ast_node_if_has_else(Node) &&
278          0 == containsLoops(isl_ast_node_if_get_else(Node), nullptr)))
279       Res = 0;
280     isl_ast_node_free(Node);
281     return Res;
282   }
283   case isl_ast_node_user:
284   default:
285     isl_ast_node_free(Node);
286     return -1;
287   }
288 }
289 
290 // Returns true when Node contains loops.
291 static bool containsLoops(__isl_take isl_ast_node *Node) {
292   return 0 == containsLoops(Node, nullptr);
293 }
294 
295 // This method is executed after the construction of a for node.
296 //
297 // It performs the following actions:
298 //
299 // - Reset the 'InParallelFor' flag, as soon as we leave a for node,
300 //   that is marked as openmp parallel.
301 //
302 static __isl_give isl_ast_node *
303 astBuildAfterFor(__isl_take isl_ast_node *Node, __isl_keep isl_ast_build *Build,
304                  void *User) {
305   isl_id *Id = isl_ast_node_get_annotation(Node);
306   if (!Id)
307     return Node;
308   struct IslAstUserPayload *Info =
309       (struct IslAstUserPayload *)isl_id_get_user(Id);
310   struct AstBuildUserInfo *BuildInfo = (struct AstBuildUserInfo *)User;
311 
312   if (Info) {
313     if (Info->IsOutermostParallel)
314       BuildInfo->InParallelFor = 0;
315     if (!containsLoops(isl_ast_node_for_get_body(Node)))
316       if (astScheduleDimIsParallel(Build, BuildInfo->Deps,
317                                    Info->IsReductionParallel))
318         Info->IsInnermostParallel = 1;
319     if (!Info->Context)
320       Info->Context = isl_ast_build_copy(Build);
321   }
322 
323   isl_id_free(Id);
324   return Node;
325 }
326 
327 static __isl_give isl_ast_node *AtEachDomain(__isl_take isl_ast_node *Node,
328                                              __isl_keep isl_ast_build *Context,
329                                              void *User) {
330   struct IslAstUserPayload *Info = nullptr;
331   isl_id *Id = isl_ast_node_get_annotation(Node);
332 
333   if (Id)
334     Info = (struct IslAstUserPayload *)isl_id_get_user(Id);
335 
336   if (!Info) {
337     // Allocate annotations once: parallel for detection might have already
338     // allocated the annotations for this node.
339     Info = allocateIslAstUser();
340     Id = isl_id_alloc(isl_ast_node_get_ctx(Node), nullptr, Info);
341     Id = isl_id_set_free_user(Id, &freeIslAstUser);
342   }
343 
344   if (!Info->Context)
345     Info->Context = isl_ast_build_copy(Context);
346 
347   return isl_ast_node_set_annotation(Node, Id);
348 }
349 
350 void IslAst::buildRunCondition(__isl_keep isl_ast_build *Context) {
351   // The conditions that need to be checked at run-time for this scop are
352   // available as an isl_set in the AssumedContext. We generate code for this
353   // check as follows. First, we generate an isl_pw_aff that is 1, if a certain
354   // combination of parameter values fulfills the conditions in the assumed
355   // context, and that is 0 otherwise. We then translate this isl_pw_aff into
356   // an isl_ast_expr. At run-time this expression can be evaluated and the
357   // optimized scop can be executed conditionally according to the result of the
358   // run-time check.
359 
360   isl_aff *Zero =
361       isl_aff_zero_on_domain(isl_local_space_from_space(S->getParamSpace()));
362   isl_aff *One =
363       isl_aff_zero_on_domain(isl_local_space_from_space(S->getParamSpace()));
364 
365   One = isl_aff_add_constant_si(One, 1);
366 
367   isl_pw_aff *PwZero = isl_pw_aff_from_aff(Zero);
368   isl_pw_aff *PwOne = isl_pw_aff_from_aff(One);
369 
370   PwOne = isl_pw_aff_intersect_domain(PwOne, S->getAssumedContext());
371   PwZero = isl_pw_aff_intersect_domain(
372       PwZero, isl_set_complement(S->getAssumedContext()));
373 
374   isl_pw_aff *Cond = isl_pw_aff_union_max(PwOne, PwZero);
375 
376   RunCondition = isl_ast_build_expr_from_pw_aff(Context, Cond);
377 }
378 
379 IslAst::IslAst(Scop *Scop, Dependences &D) : S(Scop) {
380   isl_ctx *Ctx = S->getIslCtx();
381   isl_options_set_ast_build_atomic_upper_bound(Ctx, true);
382   isl_ast_build *Context;
383   struct AstBuildUserInfo BuildInfo;
384 
385   if (UseContext)
386     Context = isl_ast_build_from_context(S->getContext());
387   else
388     Context = isl_ast_build_from_context(isl_set_universe(S->getParamSpace()));
389 
390   Context = isl_ast_build_set_at_each_domain(Context, AtEachDomain, nullptr);
391 
392   isl_union_map *Schedule =
393       isl_union_map_intersect_domain(S->getSchedule(), S->getDomains());
394 
395   Function *F = Scop->getRegion().getEntry()->getParent();
396   (void)F;
397 
398   DEBUG(dbgs() << ":: isl ast :: " << F->getName()
399                << " :: " << Scop->getRegion().getNameStr() << "\n");
400 
401   DEBUG(dbgs() << S->getContextStr() << "\n"; isl_union_map_dump(Schedule));
402 
403   if (DetectParallel || PollyVectorizerChoice != VECTORIZER_NONE) {
404     BuildInfo.Deps = &D;
405     BuildInfo.InParallelFor = 0;
406 
407     Context = isl_ast_build_set_before_each_for(Context, &astBuildBeforeFor,
408                                                 &BuildInfo);
409     Context = isl_ast_build_set_after_each_for(Context, &astBuildAfterFor,
410                                                &BuildInfo);
411   }
412 
413   buildRunCondition(Context);
414 
415   Root = isl_ast_build_ast_from_schedule(Context, Schedule);
416 
417   isl_ast_build_free(Context);
418 
419   DEBUG(pprint(dbgs()));
420 }
421 
422 IslAst::~IslAst() {
423   isl_ast_node_free(Root);
424   isl_ast_expr_free(RunCondition);
425 }
426 
427 /// Print a C like representation of the program.
428 void IslAst::pprint(llvm::raw_ostream &OS) {
429   isl_ast_node *Root;
430   isl_ast_print_options *Options;
431 
432   Options = isl_ast_print_options_alloc(S->getIslCtx());
433   Options = isl_ast_print_options_set_print_for(Options, &printFor, nullptr);
434 
435   isl_printer *P = isl_printer_to_str(S->getIslCtx());
436   P = isl_printer_set_output_format(P, ISL_FORMAT_C);
437 
438   P = isl_printer_print_ast_expr(P, RunCondition);
439   char *result = isl_printer_get_str(P);
440   P = isl_printer_flush(P);
441 
442   OS << "\nif (" << result << ")\n\n";
443   P = isl_printer_indent(P, 4);
444 
445   Root = getAst();
446   P = isl_ast_node_print(Root, P, Options);
447   result = isl_printer_get_str(P);
448   OS << result << "\n";
449   OS << "else\n";
450   OS << "    {  /* original code */ }\n\n";
451   isl_printer_free(P);
452   isl_ast_node_free(Root);
453 }
454 
455 __isl_give isl_ast_node *IslAst::getAst() { return isl_ast_node_copy(Root); }
456 __isl_give isl_ast_expr *IslAst::getRunCondition() {
457   return isl_ast_expr_copy(RunCondition);
458 }
459 
460 void IslAstInfo::pprint(llvm::raw_ostream &OS) { Ast->pprint(OS); }
461 
462 void IslAstInfo::releaseMemory() {
463   if (Ast) {
464     delete Ast;
465     Ast = 0;
466   }
467 }
468 
469 bool IslAstInfo::runOnScop(Scop &Scop) {
470   if (Ast)
471     delete Ast;
472 
473   S = &Scop;
474 
475   Dependences &D = getAnalysis<Dependences>();
476 
477   Ast = new IslAst(&Scop, D);
478 
479   return false;
480 }
481 
482 __isl_give isl_ast_node *IslAstInfo::getAst() { return Ast->getAst(); }
483 __isl_give isl_ast_expr *IslAstInfo::getRunCondition() {
484   return Ast->getRunCondition();
485 }
486 
487 void IslAstInfo::printScop(raw_ostream &OS) const {
488   Function *F = S->getRegion().getEntry()->getParent();
489 
490   OS << F->getName() << "():\n";
491 
492   Ast->pprint(OS);
493 }
494 
495 IslAstUserPayload *IslAstInfo::getNodePayload(__isl_keep isl_ast_node *Node) {
496   isl_id *Id = isl_ast_node_get_annotation(Node);
497   if (!Id)
498     return nullptr;
499   IslAstUserPayload *Payload = (IslAstUserPayload *)isl_id_get_user(Id);
500   isl_id_free(Id);
501   return Payload;
502 }
503 
504 bool IslAstInfo::isParallel(__isl_keep isl_ast_node *Node) {
505   return (isInnermostParallel(Node) || isOuterParallel(Node)) &&
506          !isReductionParallel(Node);
507 }
508 
509 bool IslAstInfo::isInnermostParallel(__isl_keep isl_ast_node *Node) {
510   IslAstUserPayload *Payload = getNodePayload(Node);
511   return Payload && Payload->IsInnermostParallel &&
512          !Payload->IsReductionParallel;
513 }
514 
515 bool IslAstInfo::isOuterParallel(__isl_keep isl_ast_node *Node) {
516   IslAstUserPayload *Payload = getNodePayload(Node);
517   return Payload && Payload->IsOutermostParallel &&
518          !Payload->IsReductionParallel;
519 }
520 
521 bool IslAstInfo::isReductionParallel(__isl_keep isl_ast_node *Node) {
522   IslAstUserPayload *Payload = getNodePayload(Node);
523   return Payload && Payload->IsReductionParallel;
524 }
525 
526 void IslAstInfo::getAnalysisUsage(AnalysisUsage &AU) const {
527   // Get the Common analysis usage of ScopPasses.
528   ScopPass::getAnalysisUsage(AU);
529   AU.addRequired<ScopInfo>();
530   AU.addRequired<Dependences>();
531 }
532 
533 char IslAstInfo::ID = 0;
534 
535 Pass *polly::createIslAstInfoPass() { return new IslAstInfo(); }
536 
537 INITIALIZE_PASS_BEGIN(IslAstInfo, "polly-ast",
538                       "Polly - Generate an AST of the SCoP (isl)", false,
539                       false);
540 INITIALIZE_PASS_DEPENDENCY(ScopInfo);
541 INITIALIZE_PASS_DEPENDENCY(Dependences);
542 INITIALIZE_PASS_END(IslAstInfo, "polly-ast",
543                     "Polly - Generate an AST from the SCoP (isl)", false, false)
544