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