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