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