Coverage Report

Created: 2018-06-19 22:08

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