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