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