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