Coverage Report

Created: 2019-02-23 12:57

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