Coverage Report

Created: 2017-08-18 19:41

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