Coverage Report

Created: 2019-04-25 15:07

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