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