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