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