Coverage Report

Created: 2017-10-03 07:32

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