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