1 //===- IslAst.cpp - isl code generator interface --------------------------===// 2 // 3 // The LLVM Compiler Infrastructure 4 // 5 // This file is distributed under the University of Illinois Open Source 6 // License. See LICENSE.TXT for details. 7 // 8 //===----------------------------------------------------------------------===// 9 // 10 // The isl code generator interface takes a Scop and generates a isl_ast. This 11 // ist_ast can either be returned directly or it can be pretty printed to 12 // stdout. 13 // 14 // A typical isl_ast output looks like this: 15 // 16 // for (c2 = max(0, ceild(n + m, 2); c2 <= min(511, floord(5 * n, 3)); c2++) { 17 // bb2(c2); 18 // } 19 // 20 //===----------------------------------------------------------------------===// 21 22 #include "polly/CodeGen/CodeGeneration.h" 23 #include "polly/CodeGen/IslAst.h" 24 #include "polly/Dependences.h" 25 #include "polly/LinkAllPasses.h" 26 #include "polly/Options.h" 27 #include "polly/ScopInfo.h" 28 #include "llvm/Support/Debug.h" 29 30 #include "isl/union_map.h" 31 #include "isl/list.h" 32 #include "isl/ast_build.h" 33 #include "isl/set.h" 34 #include "isl/map.h" 35 #include "isl/aff.h" 36 37 #define DEBUG_TYPE "polly-ast" 38 39 using namespace llvm; 40 using namespace polly; 41 42 using IslAstUserPayload = IslAstInfo::IslAstUserPayload; 43 44 static cl::opt<bool> UseContext("polly-ast-use-context", 45 cl::desc("Use context"), cl::Hidden, 46 cl::init(false), cl::ZeroOrMore, 47 cl::cat(PollyCategory)); 48 49 static cl::opt<bool> DetectParallel("polly-ast-detect-parallel", 50 cl::desc("Detect parallelism"), cl::Hidden, 51 cl::init(false), cl::ZeroOrMore, 52 cl::cat(PollyCategory)); 53 54 namespace polly { 55 class IslAst { 56 public: 57 IslAst(Scop *Scop, Dependences &D); 58 59 ~IslAst(); 60 61 /// Print a source code representation of the program. 62 void pprint(llvm::raw_ostream &OS); 63 64 __isl_give isl_ast_node *getAst(); 65 66 /// @brief Get the run-time conditions for the Scop. 67 __isl_give isl_ast_expr *getRunCondition(); 68 69 private: 70 Scop *S; 71 isl_ast_node *Root; 72 isl_ast_expr *RunCondition; 73 74 void buildRunCondition(__isl_keep isl_ast_build *Build); 75 }; 76 } // End namespace polly. 77 78 /// @brief Free an IslAstUserPayload object pointed to by @p Ptr 79 static void freeIslAstUserPayload(void *Ptr) { 80 delete ((IslAstInfo::IslAstUserPayload *)Ptr); 81 } 82 83 IslAstInfo::IslAstUserPayload::~IslAstUserPayload() { 84 isl_ast_build_free(Build); 85 isl_pw_aff_free(MinimalDependenceDistance); 86 } 87 88 /// @brief Temporary information used when building the ast. 89 struct AstBuildUserInfo { 90 /// @brief Construct and initialize the helper struct for AST creation. 91 AstBuildUserInfo() 92 : Deps(nullptr), InParallelFor(false), LastForNodeId(nullptr) {} 93 94 /// @brief The dependence information used for the parallelism check. 95 Dependences *Deps; 96 97 /// @brief Flag to indicate that we are inside a parallel for node. 98 bool InParallelFor; 99 100 /// @brief The last iterator id created for the current SCoP. 101 isl_id *LastForNodeId; 102 }; 103 104 /// @brief Print a string @p str in a single line using @p Printer. 105 static isl_printer *printLine(__isl_take isl_printer *Printer, 106 const std::string &str, 107 __isl_keep isl_pw_aff *PWA = nullptr) { 108 Printer = isl_printer_start_line(Printer); 109 Printer = isl_printer_print_str(Printer, str.c_str()); 110 if (PWA) 111 Printer = isl_printer_print_pw_aff(Printer, PWA); 112 return isl_printer_end_line(Printer); 113 } 114 115 /// @brief Return all broken reductions as a string of clauses (OpenMP style). 116 static const std::string getBrokenReductionsStr(__isl_keep isl_ast_node *Node) { 117 IslAstInfo::MemoryAccessSet *BrokenReductions; 118 std::string str; 119 120 BrokenReductions = IslAstInfo::getBrokenReductions(Node); 121 if (!BrokenReductions || BrokenReductions->empty()) 122 return ""; 123 124 // Map each type of reduction to a comma separated list of the base addresses. 125 std::map<MemoryAccess::ReductionType, std::string> Clauses; 126 for (MemoryAccess *MA : *BrokenReductions) 127 if (MA->isWrite()) 128 Clauses[MA->getReductionType()] += 129 ", " + MA->getBaseAddr()->getName().str(); 130 131 // Now print the reductions sorted by type. Each type will cause a clause 132 // like: reduction (+ : sum0, sum1, sum2) 133 for (const auto &ReductionClause : Clauses) { 134 str += " reduction ("; 135 str += MemoryAccess::getReductionOperatorStr(ReductionClause.first); 136 // Remove the first two symbols (", ") to make the output look pretty. 137 str += " : " + ReductionClause.second.substr(2) + ")"; 138 } 139 140 return str; 141 } 142 143 /// @brief Callback executed for each for node in the ast in order to print it. 144 static isl_printer *cbPrintFor(__isl_take isl_printer *Printer, 145 __isl_take isl_ast_print_options *Options, 146 __isl_keep isl_ast_node *Node, void *) { 147 148 isl_pw_aff *DD = IslAstInfo::getMinimalDependenceDistance(Node); 149 const std::string BrokenReductionsStr = getBrokenReductionsStr(Node); 150 const std::string DepDisPragmaStr = "#pragma minimal dependence distance: "; 151 const std::string SimdPragmaStr = "#pragma simd"; 152 const std::string OmpPragmaStr = "#pragma omp parallel for"; 153 154 if (DD) 155 Printer = printLine(Printer, DepDisPragmaStr, DD); 156 157 if (IslAstInfo::isInnermostParallel(Node)) 158 Printer = printLine(Printer, SimdPragmaStr + BrokenReductionsStr); 159 160 if (IslAstInfo::isOutermostParallel(Node)) 161 Printer = printLine(Printer, OmpPragmaStr + BrokenReductionsStr); 162 163 isl_pw_aff_free(DD); 164 return isl_ast_node_for_print(Node, Printer, Options); 165 } 166 167 /// @brief Check if the current scheduling dimension is parallel 168 /// 169 /// In case the dimension is parallel we also check if any reduction 170 /// dependences is broken when we exploit this parallelism. If so, 171 /// @p IsReductionParallel will be set to true. The reduction dependences we use 172 /// to check are actually the union of the transitive closure of the initial 173 /// reduction dependences together with their reveresal. Even though these 174 /// dependences connect all iterations with each other (thus they are cyclic) 175 /// we can perform the parallelism check as we are only interested in a zero 176 /// (or non-zero) dependence distance on the dimension in question. 177 static bool astScheduleDimIsParallel(__isl_keep isl_ast_build *Build, 178 Dependences *D, 179 IslAstUserPayload *NodeInfo) { 180 if (!D->hasValidDependences()) 181 return false; 182 183 isl_union_map *Schedule = isl_ast_build_get_schedule(Build); 184 isl_union_map *Deps = D->getDependences( 185 Dependences::TYPE_RAW | Dependences::TYPE_WAW | Dependences::TYPE_WAR); 186 187 if (!D->isParallel(Schedule, Deps, &NodeInfo->MinimalDependenceDistance) && 188 !isl_union_map_free(Schedule)) 189 return false; 190 191 isl_union_map *RedDeps = D->getDependences(Dependences::TYPE_TC_RED); 192 if (!D->isParallel(Schedule, RedDeps)) 193 NodeInfo->IsReductionParallel = true; 194 195 if (!NodeInfo->IsReductionParallel && !isl_union_map_free(Schedule)) 196 return true; 197 198 // Annotate reduction parallel nodes with the memory accesses which caused the 199 // reduction dependences parallel execution of the node conflicts with. 200 for (const auto &MaRedPair : D->getReductionDependences()) { 201 if (!MaRedPair.second) 202 continue; 203 RedDeps = isl_union_map_from_map(isl_map_copy(MaRedPair.second)); 204 if (!D->isParallel(Schedule, RedDeps)) 205 NodeInfo->BrokenReductions.insert(MaRedPair.first); 206 } 207 208 isl_union_map_free(Schedule); 209 return true; 210 } 211 212 // This method is executed before the construction of a for node. It creates 213 // an isl_id that is used to annotate the subsequently generated ast for nodes. 214 // 215 // In this function we also run the following analyses: 216 // 217 // - Detection of openmp parallel loops 218 // 219 static __isl_give isl_id *astBuildBeforeFor(__isl_keep isl_ast_build *Build, 220 void *User) { 221 AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User; 222 IslAstUserPayload *Payload = new IslAstUserPayload(); 223 isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", Payload); 224 Id = isl_id_set_free_user(Id, freeIslAstUserPayload); 225 BuildInfo->LastForNodeId = Id; 226 227 // Test for parallelism only if we are not already inside a parallel loop 228 if (!BuildInfo->InParallelFor) 229 BuildInfo->InParallelFor = Payload->IsOutermostParallel = 230 astScheduleDimIsParallel(Build, BuildInfo->Deps, Payload); 231 232 return Id; 233 } 234 235 // This method is executed after the construction of a for node. 236 // 237 // It performs the following actions: 238 // 239 // - Reset the 'InParallelFor' flag, as soon as we leave a for node, 240 // that is marked as openmp parallel. 241 // 242 static __isl_give isl_ast_node * 243 astBuildAfterFor(__isl_take isl_ast_node *Node, __isl_keep isl_ast_build *Build, 244 void *User) { 245 isl_id *Id = isl_ast_node_get_annotation(Node); 246 assert(Id && "Post order visit assumes annotated for nodes"); 247 IslAstUserPayload *Payload = (IslAstUserPayload *)isl_id_get_user(Id); 248 assert(Payload && "Post order visit assumes annotated for nodes"); 249 250 AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User; 251 assert(!Payload->Build && "Build environment already set"); 252 Payload->Build = isl_ast_build_copy(Build); 253 Payload->IsInnermost = (Id == BuildInfo->LastForNodeId); 254 255 // Innermost loops that are surrounded by parallel loops have not yet been 256 // tested for parallelism. Test them here to ensure we check all innermost 257 // loops for parallelism. 258 if (Payload->IsInnermost && BuildInfo->InParallelFor) { 259 if (Payload->IsOutermostParallel) 260 Payload->IsInnermostParallel = true; 261 else 262 Payload->IsInnermostParallel = 263 astScheduleDimIsParallel(Build, BuildInfo->Deps, Payload); 264 } 265 if (Payload->IsOutermostParallel) 266 BuildInfo->InParallelFor = false; 267 268 isl_id_free(Id); 269 return Node; 270 } 271 272 static __isl_give isl_ast_node *AtEachDomain(__isl_take isl_ast_node *Node, 273 __isl_keep isl_ast_build *Build, 274 void *User) { 275 assert(!isl_ast_node_get_annotation(Node) && "Node already annotated"); 276 277 IslAstUserPayload *Payload = new IslAstUserPayload(); 278 isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", Payload); 279 Id = isl_id_set_free_user(Id, freeIslAstUserPayload); 280 281 Payload->Build = isl_ast_build_copy(Build); 282 283 return isl_ast_node_set_annotation(Node, Id); 284 } 285 286 void IslAst::buildRunCondition(__isl_keep isl_ast_build *Build) { 287 // The conditions that need to be checked at run-time for this scop are 288 // available as an isl_set in the AssumedContext. We generate code for this 289 // check as follows. First, we generate an isl_pw_aff that is 1, if a certain 290 // combination of parameter values fulfills the conditions in the assumed 291 // context, and that is 0 otherwise. We then translate this isl_pw_aff into 292 // an isl_ast_expr. At run-time this expression can be evaluated and the 293 // optimized scop can be executed conditionally according to the result of the 294 // run-time check. 295 296 isl_aff *Zero = 297 isl_aff_zero_on_domain(isl_local_space_from_space(S->getParamSpace())); 298 isl_aff *One = 299 isl_aff_zero_on_domain(isl_local_space_from_space(S->getParamSpace())); 300 301 One = isl_aff_add_constant_si(One, 1); 302 303 isl_pw_aff *PwZero = isl_pw_aff_from_aff(Zero); 304 isl_pw_aff *PwOne = isl_pw_aff_from_aff(One); 305 306 PwOne = isl_pw_aff_intersect_domain(PwOne, S->getAssumedContext()); 307 PwZero = isl_pw_aff_intersect_domain( 308 PwZero, isl_set_complement(S->getAssumedContext())); 309 310 isl_pw_aff *Cond = isl_pw_aff_union_max(PwOne, PwZero); 311 312 RunCondition = isl_ast_build_expr_from_pw_aff(Build, Cond); 313 314 // Create the alias checks from the minimal/maximal accesses in each alias 315 // group. This operation is by construction quadratic in the number of 316 // elements in each alias group. 317 isl_ast_expr *NonAliasGroup, *MinExpr, *MaxExpr; 318 for (const Scop::MinMaxVectorTy *MinMaxAccesses : S->getAliasGroups()) { 319 auto AccEnd = MinMaxAccesses->end(); 320 for (auto AccIt0 = MinMaxAccesses->begin(); AccIt0 != AccEnd; ++AccIt0) { 321 for (auto AccIt1 = AccIt0 + 1; AccIt1 != AccEnd; ++AccIt1) { 322 MinExpr = 323 isl_ast_expr_address_of(isl_ast_build_access_from_pw_multi_aff( 324 Build, isl_pw_multi_aff_copy(AccIt0->first))); 325 MaxExpr = 326 isl_ast_expr_address_of(isl_ast_build_access_from_pw_multi_aff( 327 Build, isl_pw_multi_aff_copy(AccIt1->second))); 328 NonAliasGroup = isl_ast_expr_le(MaxExpr, MinExpr); 329 MinExpr = 330 isl_ast_expr_address_of(isl_ast_build_access_from_pw_multi_aff( 331 Build, isl_pw_multi_aff_copy(AccIt1->first))); 332 MaxExpr = 333 isl_ast_expr_address_of(isl_ast_build_access_from_pw_multi_aff( 334 Build, isl_pw_multi_aff_copy(AccIt0->second))); 335 NonAliasGroup = 336 isl_ast_expr_or(NonAliasGroup, isl_ast_expr_le(MaxExpr, MinExpr)); 337 RunCondition = isl_ast_expr_and(RunCondition, NonAliasGroup); 338 } 339 } 340 } 341 } 342 343 IslAst::IslAst(Scop *Scop, Dependences &D) : S(Scop) { 344 isl_ctx *Ctx = S->getIslCtx(); 345 isl_options_set_ast_build_atomic_upper_bound(Ctx, true); 346 isl_ast_build *Build; 347 AstBuildUserInfo BuildInfo; 348 349 if (UseContext) 350 Build = isl_ast_build_from_context(S->getContext()); 351 else 352 Build = isl_ast_build_from_context(isl_set_universe(S->getParamSpace())); 353 354 Build = isl_ast_build_set_at_each_domain(Build, AtEachDomain, nullptr); 355 356 isl_union_map *Schedule = 357 isl_union_map_intersect_domain(S->getSchedule(), S->getDomains()); 358 359 if (DetectParallel || PollyVectorizerChoice != VECTORIZER_NONE) { 360 BuildInfo.Deps = &D; 361 BuildInfo.InParallelFor = 0; 362 363 Build = isl_ast_build_set_before_each_for(Build, &astBuildBeforeFor, 364 &BuildInfo); 365 Build = 366 isl_ast_build_set_after_each_for(Build, &astBuildAfterFor, &BuildInfo); 367 } 368 369 buildRunCondition(Build); 370 371 Root = isl_ast_build_ast_from_schedule(Build, Schedule); 372 373 isl_ast_build_free(Build); 374 } 375 376 IslAst::~IslAst() { 377 isl_ast_node_free(Root); 378 isl_ast_expr_free(RunCondition); 379 } 380 381 __isl_give isl_ast_node *IslAst::getAst() { return isl_ast_node_copy(Root); } 382 __isl_give isl_ast_expr *IslAst::getRunCondition() { 383 return isl_ast_expr_copy(RunCondition); 384 } 385 386 void IslAstInfo::releaseMemory() { 387 if (Ast) { 388 delete Ast; 389 Ast = 0; 390 } 391 } 392 393 bool IslAstInfo::runOnScop(Scop &Scop) { 394 if (Ast) 395 delete Ast; 396 397 S = &Scop; 398 399 Dependences &D = getAnalysis<Dependences>(); 400 401 Ast = new IslAst(&Scop, D); 402 403 DEBUG(printScop(dbgs())); 404 return false; 405 } 406 407 __isl_give isl_ast_node *IslAstInfo::getAst() const { return Ast->getAst(); } 408 __isl_give isl_ast_expr *IslAstInfo::getRunCondition() const { 409 return Ast->getRunCondition(); 410 } 411 412 IslAstUserPayload *IslAstInfo::getNodePayload(__isl_keep isl_ast_node *Node) { 413 isl_id *Id = isl_ast_node_get_annotation(Node); 414 if (!Id) 415 return nullptr; 416 IslAstUserPayload *Payload = (IslAstUserPayload *)isl_id_get_user(Id); 417 isl_id_free(Id); 418 return Payload; 419 } 420 421 bool IslAstInfo::isInnermost(__isl_keep isl_ast_node *Node) { 422 IslAstUserPayload *Payload = getNodePayload(Node); 423 return Payload && Payload->IsInnermost; 424 } 425 426 bool IslAstInfo::isParallel(__isl_keep isl_ast_node *Node) { 427 return IslAstInfo::isInnermostParallel(Node) || 428 IslAstInfo::isOutermostParallel(Node); 429 } 430 431 bool IslAstInfo::isInnermostParallel(__isl_keep isl_ast_node *Node) { 432 IslAstUserPayload *Payload = getNodePayload(Node); 433 return Payload && Payload->IsInnermostParallel; 434 } 435 436 bool IslAstInfo::isOutermostParallel(__isl_keep isl_ast_node *Node) { 437 IslAstUserPayload *Payload = getNodePayload(Node); 438 return Payload && Payload->IsOutermostParallel; 439 } 440 441 bool IslAstInfo::isReductionParallel(__isl_keep isl_ast_node *Node) { 442 IslAstUserPayload *Payload = getNodePayload(Node); 443 return Payload && Payload->IsReductionParallel; 444 } 445 446 isl_union_map *IslAstInfo::getSchedule(__isl_keep isl_ast_node *Node) { 447 IslAstUserPayload *Payload = getNodePayload(Node); 448 return Payload ? isl_ast_build_get_schedule(Payload->Build) : nullptr; 449 } 450 451 isl_pw_aff * 452 IslAstInfo::getMinimalDependenceDistance(__isl_keep isl_ast_node *Node) { 453 IslAstUserPayload *Payload = getNodePayload(Node); 454 return Payload ? isl_pw_aff_copy(Payload->MinimalDependenceDistance) 455 : nullptr; 456 } 457 458 IslAstInfo::MemoryAccessSet * 459 IslAstInfo::getBrokenReductions(__isl_keep isl_ast_node *Node) { 460 IslAstUserPayload *Payload = getNodePayload(Node); 461 return Payload ? &Payload->BrokenReductions : nullptr; 462 } 463 464 isl_ast_build *IslAstInfo::getBuild(__isl_keep isl_ast_node *Node) { 465 IslAstUserPayload *Payload = getNodePayload(Node); 466 return Payload ? Payload->Build : nullptr; 467 } 468 469 void IslAstInfo::printScop(raw_ostream &OS) const { 470 isl_ast_print_options *Options; 471 isl_ast_node *RootNode = getAst(); 472 isl_ast_expr *RunCondition = getRunCondition(); 473 char *RtCStr, *AstStr; 474 475 Scop &S = getCurScop(); 476 Options = isl_ast_print_options_alloc(S.getIslCtx()); 477 Options = isl_ast_print_options_set_print_for(Options, cbPrintFor, nullptr); 478 479 isl_printer *P = isl_printer_to_str(S.getIslCtx()); 480 P = isl_printer_print_ast_expr(P, RunCondition); 481 RtCStr = isl_printer_get_str(P); 482 P = isl_printer_flush(P); 483 P = isl_printer_indent(P, 4); 484 P = isl_printer_set_output_format(P, ISL_FORMAT_C); 485 P = isl_ast_node_print(RootNode, P, Options); 486 AstStr = isl_printer_get_str(P); 487 488 Function *F = S.getRegion().getEntry()->getParent(); 489 isl_union_map *Schedule = 490 isl_union_map_intersect_domain(S.getSchedule(), S.getDomains()); 491 492 OS << ":: isl ast :: " << F->getName() << " :: " << S.getRegion().getNameStr() 493 << "\n"; 494 DEBUG(dbgs() << S.getContextStr() << "\n"; isl_union_map_dump(Schedule)); 495 OS << "\nif (" << RtCStr << ")\n\n"; 496 OS << AstStr << "\n"; 497 OS << "else\n"; 498 OS << " { /* original code */ }\n\n"; 499 500 isl_ast_expr_free(RunCondition); 501 isl_union_map_free(Schedule); 502 isl_ast_node_free(RootNode); 503 isl_printer_free(P); 504 } 505 506 void IslAstInfo::getAnalysisUsage(AnalysisUsage &AU) const { 507 // Get the Common analysis usage of ScopPasses. 508 ScopPass::getAnalysisUsage(AU); 509 AU.addRequired<ScopInfo>(); 510 AU.addRequired<Dependences>(); 511 } 512 513 char IslAstInfo::ID = 0; 514 515 Pass *polly::createIslAstInfoPass() { return new IslAstInfo(); } 516 517 INITIALIZE_PASS_BEGIN(IslAstInfo, "polly-ast", 518 "Polly - Generate an AST of the SCoP (isl)", false, 519 false); 520 INITIALIZE_PASS_DEPENDENCY(ScopInfo); 521 INITIALIZE_PASS_DEPENDENCY(Dependences); 522 INITIALIZE_PASS_END(IslAstInfo, "polly-ast", 523 "Polly - Generate an AST from the SCoP (isl)", false, false) 524