Coverage Report

Created: 2017-06-28 17:40

/Users/buildslave/jenkins/sharedspace/clang-stage2-coverage-R@2/llvm/tools/polly/lib/CodeGen/IslAst.cpp
Line
Count
Source (jump to first uncovered line)
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 an 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
// An in-depth discussion of our AST generation approach can be found in:
21
//
22
// Polyhedral AST generation is more than scanning polyhedra
23
// Tobias Grosser, Sven Verdoolaege, Albert Cohen
24
// ACM Transactions on Programming Languages and Systems (TOPLAS),
25
// 37(4), July 2015
26
// http://www.grosser.es/#pub-polyhedral-AST-generation
27
//
28
//===----------------------------------------------------------------------===//
29
30
#include "polly/CodeGen/IslAst.h"
31
#include "polly/CodeGen/CodeGeneration.h"
32
#include "polly/DependenceInfo.h"
33
#include "polly/LinkAllPasses.h"
34
#include "polly/Options.h"
35
#include "polly/ScopInfo.h"
36
#include "polly/Support/GICHelper.h"
37
#include "llvm/Analysis/RegionInfo.h"
38
#include "llvm/Support/Debug.h"
39
#include "isl/aff.h"
40
#include "isl/ast_build.h"
41
#include "isl/list.h"
42
#include "isl/map.h"
43
#include "isl/set.h"
44
#include "isl/union_map.h"
45
46
#define DEBUG_TYPE "polly-ast"
47
48
using namespace llvm;
49
using namespace polly;
50
51
using IslAstUserPayload = IslAstInfo::IslAstUserPayload;
52
53
static cl::opt<bool>
54
    PollyParallel("polly-parallel",
55
                  cl::desc("Generate thread parallel code (isl codegen only)"),
56
                  cl::init(false), cl::ZeroOrMore, cl::cat(PollyCategory));
57
58
static cl::opt<bool> PollyParallelForce(
59
    "polly-parallel-force",
60
    cl::desc(
61
        "Force generation of thread parallel code ignoring any cost model"),
62
    cl::init(false), cl::ZeroOrMore, cl::cat(PollyCategory));
63
64
static cl::opt<bool> UseContext("polly-ast-use-context",
65
                                cl::desc("Use context"), cl::Hidden,
66
                                cl::init(true), cl::ZeroOrMore,
67
                                cl::cat(PollyCategory));
68
69
static cl::opt<bool> DetectParallel("polly-ast-detect-parallel",
70
                                    cl::desc("Detect parallelism"), cl::Hidden,
71
                                    cl::init(false), cl::ZeroOrMore,
72
                                    cl::cat(PollyCategory));
73
74
namespace polly {
75
/// Temporary information used when building the ast.
76
struct AstBuildUserInfo {
77
  /// Construct and initialize the helper struct for AST creation.
78
  AstBuildUserInfo()
79
444
      : Deps(nullptr), InParallelFor(false), LastForNodeId(nullptr) {}
80
81
  /// The dependence information used for the parallelism check.
82
  const Dependences *Deps;
83
84
  /// Flag to indicate that we are inside a parallel for node.
85
  bool InParallelFor;
86
87
  /// The last iterator id created for the current SCoP.
88
  isl_id *LastForNodeId;
89
};
90
} // namespace polly
91
92
/// Free an IslAstUserPayload object pointed to by @p Ptr.
93
1.56k
static void freeIslAstUserPayload(void *Ptr) {
94
1.56k
  delete ((IslAstInfo::IslAstUserPayload *)Ptr);
95
1.56k
}
96
97
1.56k
IslAstInfo::IslAstUserPayload::~IslAstUserPayload() {
98
1.56k
  isl_ast_build_free(Build);
99
1.56k
  isl_pw_aff_free(MinimalDependenceDistance);
100
1.56k
}
101
102
/// Print a string @p str in a single line using @p Printer.
103
static isl_printer *printLine(__isl_take isl_printer *Printer,
104
                              const std::string &str,
105
120
                              __isl_keep isl_pw_aff *PWA = nullptr) {
106
120
  Printer = isl_printer_start_line(Printer);
107
120
  Printer = isl_printer_print_str(Printer, str.c_str());
108
120
  if (PWA)
109
21
    Printer = isl_printer_print_pw_aff(Printer, PWA);
110
120
  return isl_printer_end_line(Printer);
111
120
}
112
113
/// Return all broken reductions as a string of clauses (OpenMP style).
114
370
static const std::string getBrokenReductionsStr(__isl_keep isl_ast_node *Node) {
115
370
  IslAstInfo::MemoryAccessSet *BrokenReductions;
116
370
  std::string str;
117
370
118
370
  BrokenReductions = IslAstInfo::getBrokenReductions(Node);
119
370
  if (
!BrokenReductions || 370
BrokenReductions->empty()154
)
120
347
    return "";
121
370
122
370
  // Map each type of reduction to a comma separated list of the base addresses.
123
23
  std::map<MemoryAccess::ReductionType, std::string> Clauses;
124
23
  for (MemoryAccess *MA : *BrokenReductions)
125
54
    
if (54
MA->isWrite()54
)
126
27
      Clauses[MA->getReductionType()] +=
127
27
          ", " + MA->getScopArrayInfo()->getName();
128
23
129
23
  // Now print the reductions sorted by type. Each type will cause a clause
130
23
  // like:  reduction (+ : sum0, sum1, sum2)
131
26
  for (const auto &ReductionClause : Clauses) {
132
26
    str += " reduction (";
133
26
    str += MemoryAccess::getReductionOperatorStr(ReductionClause.first);
134
26
    // Remove the first two symbols (", ") to make the output look pretty.
135
26
    str += " : " + ReductionClause.second.substr(2) + ")";
136
26
  }
137
23
138
23
  return str;
139
370
}
140
141
/// Callback executed for each for node in the ast in order to print it.
142
static isl_printer *cbPrintFor(__isl_take isl_printer *Printer,
143
                               __isl_take isl_ast_print_options *Options,
144
370
                               __isl_keep isl_ast_node *Node, void *) {
145
370
146
370
  isl_pw_aff *DD = IslAstInfo::getMinimalDependenceDistance(Node);
147
370
  const std::string BrokenReductionsStr = getBrokenReductionsStr(Node);
148
370
  const std::string KnownParallelStr = "#pragma known-parallel";
149
370
  const std::string DepDisPragmaStr = "#pragma minimal dependence distance: ";
150
370
  const std::string SimdPragmaStr = "#pragma simd";
151
370
  const std::string OmpPragmaStr = "#pragma omp parallel for";
152
370
153
370
  if (DD)
154
21
    Printer = printLine(Printer, DepDisPragmaStr, DD);
155
370
156
370
  if (IslAstInfo::isInnermostParallel(Node))
157
46
    Printer = printLine(Printer, SimdPragmaStr + BrokenReductionsStr);
158
370
159
370
  if (IslAstInfo::isExecutedInParallel(Node))
160
18
    Printer = printLine(Printer, OmpPragmaStr);
161
352
  else 
if (352
IslAstInfo::isOutermostParallel(Node)352
)
162
35
    Printer = printLine(Printer, KnownParallelStr + BrokenReductionsStr);
163
370
164
370
  isl_pw_aff_free(DD);
165
370
  return isl_ast_node_for_print(Node, Printer, Options);
166
370
}
167
168
/// Check if the current scheduling dimension is parallel.
169
///
170
/// In case the dimension is parallel we also check if any reduction
171
/// dependences is broken when we exploit this parallelism. If so,
172
/// @p IsReductionParallel will be set to true. The reduction dependences we use
173
/// to check are actually the union of the transitive closure of the initial
174
/// reduction dependences together with their reversal. Even though these
175
/// dependences connect all iterations with each other (thus they are cyclic)
176
/// we can perform the parallelism check as we are only interested in a zero
177
/// (or non-zero) dependence distance on the dimension in question.
178
static bool astScheduleDimIsParallel(__isl_keep isl_ast_build *Build,
179
                                     const Dependences *D,
180
196
                                     IslAstUserPayload *NodeInfo) {
181
196
  if (!D->hasValidDependences())
182
1
    return false;
183
196
184
195
  isl_union_map *Schedule = isl_ast_build_get_schedule(Build);
185
195
  isl_union_map *Deps = D->getDependences(
186
195
      Dependences::TYPE_RAW | Dependences::TYPE_WAW | Dependences::TYPE_WAR);
187
195
188
195
  if (!D->isParallel(Schedule, Deps, &NodeInfo->MinimalDependenceDistance) &&
189
29
      !isl_union_map_free(Schedule))
190
29
    return false;
191
195
192
166
  isl_union_map *RedDeps = D->getDependences(Dependences::TYPE_TC_RED);
193
166
  if (!D->isParallel(Schedule, RedDeps))
194
38
    NodeInfo->IsReductionParallel = true;
195
166
196
166
  if (
!NodeInfo->IsReductionParallel && 166
!isl_union_map_free(Schedule)128
)
197
128
    return true;
198
166
199
166
  // Annotate reduction parallel nodes with the memory accesses which caused the
200
166
  // reduction dependences parallel execution of the node conflicts with.
201
84
  
for (const auto &MaRedPair : D->getReductionDependences()) 38
{84
202
84
    if (!MaRedPair.second)
203
0
      continue;
204
84
    RedDeps = isl_union_map_from_map(isl_map_copy(MaRedPair.second));
205
84
    if (!D->isParallel(Schedule, RedDeps))
206
84
      NodeInfo->BrokenReductions.insert(MaRedPair.first);
207
84
  }
208
38
209
38
  isl_union_map_free(Schedule);
210
38
  return true;
211
166
}
212
213
// This method is executed before the construction of a for node. It creates
214
// an isl_id that is used to annotate the subsequently generated ast for nodes.
215
//
216
// In this function we also run the following analyses:
217
//
218
// - Detection of openmp parallel loops
219
//
220
static __isl_give isl_id *astBuildBeforeFor(__isl_keep isl_ast_build *Build,
221
258
                                            void *User) {
222
258
  AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
223
258
  IslAstUserPayload *Payload = new IslAstUserPayload();
224
258
  isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", Payload);
225
258
  Id = isl_id_set_free_user(Id, freeIslAstUserPayload);
226
258
  BuildInfo->LastForNodeId = Id;
227
258
228
258
  // Test for parallelism only if we are not already inside a parallel loop
229
258
  if (!BuildInfo->InParallelFor)
230
161
    BuildInfo->InParallelFor = Payload->IsOutermostParallel =
231
161
        astScheduleDimIsParallel(Build, BuildInfo->Deps, Payload);
232
258
233
258
  return Id;
234
258
}
235
236
// This method is executed after the construction of a for node.
237
//
238
// It performs the following actions:
239
//
240
// - Reset the 'InParallelFor' flag, as soon as we leave a for node,
241
//   that is marked as openmp parallel.
242
//
243
static __isl_give isl_ast_node *
244
astBuildAfterFor(__isl_take isl_ast_node *Node, __isl_keep isl_ast_build *Build,
245
258
                 void *User) {
246
258
  isl_id *Id = isl_ast_node_get_annotation(Node);
247
258
  assert(Id && "Post order visit assumes annotated for nodes");
248
258
  IslAstUserPayload *Payload = (IslAstUserPayload *)isl_id_get_user(Id);
249
258
  assert(Payload && "Post order visit assumes annotated for nodes");
250
258
251
258
  AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
252
258
  assert(!Payload->Build && "Build environment already set");
253
258
  Payload->Build = isl_ast_build_copy(Build);
254
258
  Payload->IsInnermost = (Id == BuildInfo->LastForNodeId);
255
258
256
258
  // Innermost loops that are surrounded by parallel loops have not yet been
257
258
  // tested for parallelism. Test them here to ensure we check all innermost
258
258
  // loops for parallelism.
259
258
  if (
Payload->IsInnermost && 258
BuildInfo->InParallelFor156
)
{143
260
143
    if (
Payload->IsOutermostParallel143
)
{94
261
94
      Payload->IsInnermostParallel = true;
262
49
    } else {
263
49
      if (PollyVectorizerChoice == VECTORIZER_NONE)
264
35
        Payload->IsInnermostParallel =
265
35
            astScheduleDimIsParallel(Build, BuildInfo->Deps, Payload);
266
49
    }
267
143
  }
268
258
  if (Payload->IsOutermostParallel)
269
136
    BuildInfo->InParallelFor = false;
270
258
271
258
  isl_id_free(Id);
272
258
  return Node;
273
258
}
274
275
static isl_stat astBuildBeforeMark(__isl_keep isl_id *MarkId,
276
                                   __isl_keep isl_ast_build *Build,
277
40
                                   void *User) {
278
40
  if (!MarkId)
279
0
    return isl_stat_error;
280
40
281
40
  AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
282
40
  if (!strcmp(isl_id_get_name(MarkId), "SIMD"))
283
14
    BuildInfo->InParallelFor = true;
284
40
285
40
  return isl_stat_ok;
286
40
}
287
288
static __isl_give isl_ast_node *
289
astBuildAfterMark(__isl_take isl_ast_node *Node,
290
40
                  __isl_keep isl_ast_build *Build, void *User) {
291
40
  assert(isl_ast_node_get_type(Node) == isl_ast_node_mark);
292
40
  AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
293
40
  auto *Id = isl_ast_node_mark_get_id(Node);
294
40
  if (!strcmp(isl_id_get_name(Id), "SIMD"))
295
14
    BuildInfo->InParallelFor = false;
296
40
  isl_id_free(Id);
297
40
  return Node;
298
40
}
299
300
static __isl_give isl_ast_node *AtEachDomain(__isl_take isl_ast_node *Node,
301
                                             __isl_keep isl_ast_build *Build,
302
1.31k
                                             void *User) {
303
1.31k
  assert(!isl_ast_node_get_annotation(Node) && "Node already annotated");
304
1.31k
305
1.31k
  IslAstUserPayload *Payload = new IslAstUserPayload();
306
1.31k
  isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", Payload);
307
1.31k
  Id = isl_id_set_free_user(Id, freeIslAstUserPayload);
308
1.31k
309
1.31k
  Payload->Build = isl_ast_build_copy(Build);
310
1.31k
311
1.31k
  return isl_ast_node_set_annotation(Node, Id);
312
1.31k
}
313
314
// Build alias check condition given a pair of minimal/maximal access.
315
static __isl_give isl_ast_expr *
316
buildCondition(__isl_keep isl_ast_build *Build, const Scop::MinMaxAccessTy *It0,
317
155
               const Scop::MinMaxAccessTy *It1) {
318
155
  isl_ast_expr *NonAliasGroup, *MinExpr, *MaxExpr;
319
155
  MinExpr = isl_ast_expr_address_of(isl_ast_build_access_from_pw_multi_aff(
320
155
      Build, isl_pw_multi_aff_copy(It0->first)));
321
155
  MaxExpr = isl_ast_expr_address_of(isl_ast_build_access_from_pw_multi_aff(
322
155
      Build, isl_pw_multi_aff_copy(It1->second)));
323
155
  NonAliasGroup = isl_ast_expr_le(MaxExpr, MinExpr);
324
155
  MinExpr = isl_ast_expr_address_of(isl_ast_build_access_from_pw_multi_aff(
325
155
      Build, isl_pw_multi_aff_copy(It1->first)));
326
155
  MaxExpr = isl_ast_expr_address_of(isl_ast_build_access_from_pw_multi_aff(
327
155
      Build, isl_pw_multi_aff_copy(It0->second)));
328
155
  NonAliasGroup =
329
155
      isl_ast_expr_or(NonAliasGroup, isl_ast_expr_le(MaxExpr, MinExpr));
330
155
331
155
  return NonAliasGroup;
332
155
}
333
334
__isl_give isl_ast_expr *
335
444
IslAst::buildRunCondition(Scop &S, __isl_keep isl_ast_build *Build) {
336
444
  isl_ast_expr *RunCondition;
337
444
338
444
  // The conditions that need to be checked at run-time for this scop are
339
444
  // available as an isl_set in the runtime check context from which we can
340
444
  // directly derive a run-time condition.
341
444
  auto *PosCond = isl_ast_build_expr_from_set(Build, S.getAssumedContext());
342
444
  if (
S.hasTrivialInvalidContext()444
)
{312
343
312
    RunCondition = PosCond;
344
132
  } else {
345
132
    auto *ZeroV = isl_val_zero(isl_ast_build_get_ctx(Build));
346
132
    auto *NegCond = isl_ast_build_expr_from_set(Build, S.getInvalidContext());
347
132
    auto *NotNegCond = isl_ast_expr_eq(isl_ast_expr_from_val(ZeroV), NegCond);
348
132
    RunCondition = isl_ast_expr_and(PosCond, NotNegCond);
349
132
  }
350
444
351
444
  // Create the alias checks from the minimal/maximal accesses in each alias
352
444
  // group which consists of read only and non read only (read write) accesses.
353
444
  // This operation is by construction quadratic in the read-write pointers and
354
444
  // linear in the read only pointers in each alias group.
355
96
  for (const Scop::MinMaxVectorPairTy &MinMaxAccessPair : S.getAliasGroups()) {
356
96
    auto &MinMaxReadWrite = MinMaxAccessPair.first;
357
96
    auto &MinMaxReadOnly = MinMaxAccessPair.second;
358
96
    auto RWAccEnd = MinMaxReadWrite.end();
359
96
360
206
    for (auto RWAccIt0 = MinMaxReadWrite.begin(); RWAccIt0 != RWAccEnd;
361
110
         
++RWAccIt0110
)
{110
362
125
      for (auto RWAccIt1 = RWAccIt0 + 1; 
RWAccIt1 != RWAccEnd125
;
++RWAccIt115
)
363
15
        RunCondition = isl_ast_expr_and(
364
15
            RunCondition, buildCondition(Build, RWAccIt0, RWAccIt1));
365
110
      for (const Scop::MinMaxAccessTy &ROAccIt : MinMaxReadOnly)
366
140
        RunCondition = isl_ast_expr_and(
367
140
            RunCondition, buildCondition(Build, RWAccIt0, &ROAccIt));
368
110
    }
369
96
  }
370
444
371
444
  return RunCondition;
372
444
}
373
374
/// Simple cost analysis for a given SCoP.
375
///
376
/// TODO: Improve this analysis and extract it to make it usable in other
377
///       places too.
378
///       In order to improve the cost model we could either keep track of
379
///       performed optimizations (e.g., tiling) or compute properties on the
380
///       original as well as optimized SCoP (e.g., #stride-one-accesses).
381
444
static bool benefitsFromPolly(Scop &Scop, bool PerformParallelTest) {
382
444
383
444
  if (PollyProcessUnprofitable)
384
444
    return true;
385
444
386
444
  // Check if nothing interesting happened.
387
0
  
if (0
!PerformParallelTest && 0
!Scop.isOptimized()0
&&
388
0
      Scop.getAliasGroups().empty())
389
0
    return false;
390
0
391
0
  // The default assumption is that Polly improves the code.
392
0
  return true;
393
0
}
394
395
IslAst::IslAst(Scop &Scop)
396
    : S(Scop), Root(nullptr), RunCondition(nullptr),
397
444
      Ctx(Scop.getSharedIslCtx()) {}
398
399
444
void IslAst::init(const Dependences &D) {
400
398
  bool PerformParallelTest = PollyParallel || DetectParallel ||
401
350
                             PollyVectorizerChoice != VECTORIZER_NONE;
402
444
403
444
  // We can not perform the dependence analysis and, consequently,
404
444
  // the parallel code generation in case the schedule tree contains
405
444
  // extension nodes.
406
444
  auto *ScheduleTree = S.getScheduleTree();
407
444
  PerformParallelTest =
408
124
      PerformParallelTest && !S.containsExtensionNode(ScheduleTree);
409
444
  isl_schedule_free(ScheduleTree);
410
444
411
444
  // Skip AST and code generation if there was no benefit achieved.
412
444
  if (!benefitsFromPolly(S, PerformParallelTest))
413
0
    return;
414
444
415
444
  isl_ctx *Ctx = S.getIslCtx();
416
444
  isl_options_set_ast_build_atomic_upper_bound(Ctx, true);
417
444
  isl_options_set_ast_build_detect_min_max(Ctx, true);
418
444
  isl_ast_build *Build;
419
444
  AstBuildUserInfo BuildInfo;
420
444
421
444
  if (UseContext)
422
444
    Build = isl_ast_build_from_context(S.getContext());
423
444
  else
424
0
    Build = isl_ast_build_from_context(isl_set_universe(S.getParamSpace()));
425
444
426
444
  Build = isl_ast_build_set_at_each_domain(Build, AtEachDomain, nullptr);
427
444
428
444
  if (
PerformParallelTest444
)
{124
429
124
    BuildInfo.Deps = &D;
430
124
    BuildInfo.InParallelFor = 0;
431
124
432
124
    Build = isl_ast_build_set_before_each_for(Build, &astBuildBeforeFor,
433
124
                                              &BuildInfo);
434
124
    Build =
435
124
        isl_ast_build_set_after_each_for(Build, &astBuildAfterFor, &BuildInfo);
436
124
437
124
    Build = isl_ast_build_set_before_each_mark(Build, &astBuildBeforeMark,
438
124
                                               &BuildInfo);
439
124
440
124
    Build = isl_ast_build_set_after_each_mark(Build, &astBuildAfterMark,
441
124
                                              &BuildInfo);
442
124
  }
443
444
444
444
  RunCondition = buildRunCondition(S, Build);
445
444
446
444
  Root = isl_ast_build_node_from_schedule(Build, S.getScheduleTree());
447
444
448
444
  isl_ast_build_free(Build);
449
444
}
450
451
444
IslAst IslAst::create(Scop &Scop, const Dependences &D) {
452
444
  IslAst Ast{Scop};
453
444
  Ast.init(D);
454
444
  return Ast;
455
444
}
456
457
IslAst::IslAst(IslAst &&O)
458
0
    : S(O.S), Root(O.Root), RunCondition(O.RunCondition), Ctx(O.Ctx) {
459
0
  O.Root = nullptr;
460
0
  O.RunCondition = nullptr;
461
0
}
462
463
444
IslAst::~IslAst() {
464
444
  isl_ast_node_free(Root);
465
444
  isl_ast_expr_free(RunCondition);
466
444
}
467
468
429
__isl_give isl_ast_node *IslAst::getAst() { return isl_ast_node_copy(Root); }
469
424
__isl_give isl_ast_expr *IslAst::getRunCondition() {
470
424
  return isl_ast_expr_copy(RunCondition);
471
424
}
472
473
280
__isl_give isl_ast_node *IslAstInfo::getAst() { return Ast.getAst(); }
474
275
__isl_give isl_ast_expr *IslAstInfo::getRunCondition() {
475
275
  return Ast.getRunCondition();
476
275
}
477
478
2.53k
IslAstUserPayload *IslAstInfo::getNodePayload(__isl_keep isl_ast_node *Node) {
479
2.53k
  isl_id *Id = isl_ast_node_get_annotation(Node);
480
2.53k
  if (!Id)
481
1.26k
    return nullptr;
482
1.27k
  IslAstUserPayload *Payload = (IslAstUserPayload *)isl_id_get_user(Id);
483
1.27k
  isl_id_free(Id);
484
1.27k
  return Payload;
485
2.53k
}
486
487
30
bool IslAstInfo::isInnermost(__isl_keep isl_ast_node *Node) {
488
30
  IslAstUserPayload *Payload = getNodePayload(Node);
489
30
  return Payload && Payload->IsInnermost;
490
30
}
491
492
227
bool IslAstInfo::isParallel(__isl_keep isl_ast_node *Node) {
493
227
  return IslAstInfo::isInnermostParallel(Node) ||
494
214
         IslAstInfo::isOutermostParallel(Node);
495
227
}
496
497
624
bool IslAstInfo::isInnermostParallel(__isl_keep isl_ast_node *Node) {
498
624
  IslAstUserPayload *Payload = getNodePayload(Node);
499
210
  return Payload && Payload->IsInnermostParallel;
500
624
}
501
502
622
bool IslAstInfo::isOutermostParallel(__isl_keep isl_ast_node *Node) {
503
622
  IslAstUserPayload *Payload = getNodePayload(Node);
504
208
  return Payload && Payload->IsOutermostParallel;
505
622
}
506
507
80
bool IslAstInfo::isReductionParallel(__isl_keep isl_ast_node *Node) {
508
80
  IslAstUserPayload *Payload = getNodePayload(Node);
509
80
  return Payload && Payload->IsReductionParallel;
510
80
}
511
512
623
bool IslAstInfo::isExecutedInParallel(__isl_keep isl_ast_node *Node) {
513
623
514
623
  if (!PollyParallel)
515
553
    return false;
516
623
517
623
  // Do not parallelize innermost loops.
518
623
  //
519
623
  // Parallelizing innermost loops is often not profitable, especially if
520
623
  // they have a low number of iterations.
521
623
  //
522
623
  // TODO: Decide this based on the number of loop iterations that will be
523
623
  //       executed. This can possibly require run-time checks, which again
524
623
  //       raises the question of both run-time check overhead and code size
525
623
  //       costs.
526
70
  
if (70
!PollyParallelForce && 70
isInnermost(Node)30
)
527
14
    return false;
528
70
529
56
  
return isOutermostParallel(Node) && 56
!isReductionParallel(Node)44
;
530
70
}
531
532
__isl_give isl_union_map *
533
45
IslAstInfo::getSchedule(__isl_keep isl_ast_node *Node) {
534
45
  IslAstUserPayload *Payload = getNodePayload(Node);
535
45
  return Payload ? 
isl_ast_build_get_schedule(Payload->Build)45
:
nullptr0
;
536
45
}
537
538
__isl_give isl_pw_aff *
539
370
IslAstInfo::getMinimalDependenceDistance(__isl_keep isl_ast_node *Node) {
540
370
  IslAstUserPayload *Payload = getNodePayload(Node);
541
154
  return Payload ? isl_pw_aff_copy(Payload->MinimalDependenceDistance)
542
216
                 : nullptr;
543
370
}
544
545
IslAstInfo::MemoryAccessSet *
546
370
IslAstInfo::getBrokenReductions(__isl_keep isl_ast_node *Node) {
547
370
  IslAstUserPayload *Payload = getNodePayload(Node);
548
216
  return Payload ? 
&Payload->BrokenReductions154
:
nullptr216
;
549
370
}
550
551
396
isl_ast_build *IslAstInfo::getBuild(__isl_keep isl_ast_node *Node) {
552
396
  IslAstUserPayload *Payload = getNodePayload(Node);
553
396
  return Payload ? 
Payload->Build396
:
nullptr0
;
554
396
}
555
556
IslAstInfo IslAstAnalysis::run(Scop &S, ScopAnalysisManager &SAM,
557
0
                               ScopStandardAnalysisResults &SAR) {
558
0
  return {S, SAM.getResult<DependenceAnalysis>(S, SAR).getDependences(
559
0
                 Dependences::AL_Statement)};
560
0
}
561
562
149
void IslAstInfo::print(raw_ostream &OS) {
563
149
  isl_ast_print_options *Options;
564
149
  isl_ast_node *RootNode = Ast.getAst();
565
149
  Function &F = S.getFunction();
566
149
567
149
  OS << ":: isl ast :: " << F.getName() << " :: " << S.getNameStr() << "\n";
568
149
569
149
  if (
!RootNode149
)
{0
570
0
    OS << ":: isl ast generation and code generation was skipped!\n\n";
571
0
    OS << ":: This is either because no useful optimizations could be applied "
572
0
          "(use -polly-process-unprofitable to enforce code generation) or "
573
0
          "because earlier passes such as dependence analysis timed out (use "
574
0
          "-polly-dependences-computeout=0 to set dependence analysis timeout "
575
0
          "to infinity)\n\n";
576
0
    return;
577
0
  }
578
149
579
149
  isl_ast_expr *RunCondition = Ast.getRunCondition();
580
149
  char *RtCStr, *AstStr;
581
149
582
149
  Options = isl_ast_print_options_alloc(S.getIslCtx());
583
149
  Options = isl_ast_print_options_set_print_for(Options, cbPrintFor, nullptr);
584
149
585
149
  isl_printer *P = isl_printer_to_str(S.getIslCtx());
586
149
  P = isl_printer_set_output_format(P, ISL_FORMAT_C);
587
149
  P = isl_printer_print_ast_expr(P, RunCondition);
588
149
  RtCStr = isl_printer_get_str(P);
589
149
  P = isl_printer_flush(P);
590
149
  P = isl_printer_indent(P, 4);
591
149
  P = isl_ast_node_print(RootNode, P, Options);
592
149
  AstStr = isl_printer_get_str(P);
593
149
594
149
  auto *Schedule = S.getScheduleTree();
595
149
596
149
  DEBUG({
597
149
    dbgs() << S.getContextStr() << "\n";
598
149
    dbgs() << stringFromIslObj(Schedule);
599
149
  });
600
149
  OS << "\nif (" << RtCStr << ")\n\n";
601
149
  OS << AstStr << "\n";
602
149
  OS << "else\n";
603
149
  OS << "    {  /* original code */ }\n\n";
604
149
605
149
  free(RtCStr);
606
149
  free(AstStr);
607
149
608
149
  isl_ast_expr_free(RunCondition);
609
149
  isl_schedule_free(Schedule);
610
149
  isl_ast_node_free(RootNode);
611
149
  isl_printer_free(P);
612
149
}
613
614
AnalysisKey IslAstAnalysis::Key;
615
PreservedAnalyses IslAstPrinterPass::run(Scop &S, ScopAnalysisManager &SAM,
616
                                         ScopStandardAnalysisResults &SAR,
617
0
                                         SPMUpdater &U) {
618
0
619
0
  auto &Ast = SAM.getResult<IslAstAnalysis>(S, SAR);
620
0
  Ast.print(Stream);
621
0
  return PreservedAnalyses::all();
622
0
}
623
624
1.71k
void IslAstInfoWrapperPass::releaseMemory() { Ast.reset(); }
625
626
444
bool IslAstInfoWrapperPass::runOnScop(Scop &Scop) {
627
444
  const Dependences &D =
628
444
      getAnalysis<DependenceInfo>().getDependences(Dependences::AL_Statement);
629
444
630
444
  Ast.reset(new IslAstInfo(Scop, D));
631
444
632
444
  DEBUG(printScop(dbgs(), Scop));
633
444
  return false;
634
444
}
635
466
void IslAstInfoWrapperPass::getAnalysisUsage(AnalysisUsage &AU) const {
636
466
  // Get the Common analysis usage of ScopPasses.
637
466
  ScopPass::getAnalysisUsage(AU);
638
466
  AU.addRequired<ScopInfoRegionPass>();
639
466
  AU.addRequired<DependenceInfo>();
640
466
}
641
642
149
void IslAstInfoWrapperPass::printScop(raw_ostream &OS, Scop &S) const {
643
149
  if (Ast)
644
149
    Ast->print(OS);
645
149
}
646
647
char IslAstInfoWrapperPass::ID = 0;
648
649
0
Pass *polly::createIslAstInfoWrapperPassPass() {
650
0
  return new IslAstInfoWrapperPass();
651
0
}
652
653
41.0k
INITIALIZE_PASS_BEGIN41.0k
(IslAstInfoWrapperPass, "polly-ast",41.0k
654
41.0k
                      "Polly - Generate an AST of the SCoP (isl)", false,
655
41.0k
                      false);
656
41.0k
INITIALIZE_PASS_DEPENDENCY(ScopInfoRegionPass);
657
41.0k
INITIALIZE_PASS_DEPENDENCY(DependenceInfo);
658
41.0k
INITIALIZE_PASS_END(IslAstInfoWrapperPass, "polly-ast",
659
                    "Polly - Generate an AST from the SCoP (isl)", false, false)