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 } 86 87 /// @brief Temporary information used when building the ast. 88 struct AstBuildUserInfo { 89 /// @brief Construct and initialize the helper struct for AST creation. 90 AstBuildUserInfo() 91 : Deps(nullptr), InParallelFor(false), LastForNodeId(nullptr) {} 92 93 /// @brief The dependence information used for the parallelism check. 94 Dependences *Deps; 95 96 /// @brief Flag to indicate that we are inside a parallel for node. 97 bool InParallelFor; 98 99 /// @brief The last iterator id created for the current SCoP. 100 isl_id *LastForNodeId; 101 }; 102 103 /// @brief Print a string @p str in a single line using @p Printer. 104 static isl_printer *printLine(__isl_take isl_printer *Printer, 105 const std::string &str) { 106 Printer = isl_printer_start_line(Printer); 107 Printer = isl_printer_print_str(Printer, str.c_str()); 108 return isl_printer_end_line(Printer); 109 } 110 111 /// @brief Return all broken reductions as a string of clauses (OpenMP style). 112 static const std::string getBrokenReductionsStr(__isl_keep isl_ast_node *Node) { 113 IslAstInfo::MemoryAccessSet *BrokenReductions; 114 std::string str; 115 116 BrokenReductions = IslAstInfo::getBrokenReductions(Node); 117 if (!BrokenReductions || BrokenReductions->empty()) 118 return ""; 119 120 // Map each type of reduction to a comma separated list of the base addresses. 121 std::map<MemoryAccess::ReductionType, std::string> Clauses; 122 for (MemoryAccess *MA : *BrokenReductions) 123 if (MA->isWrite()) 124 Clauses[MA->getReductionType()] += 125 ", " + MA->getBaseAddr()->getName().str(); 126 127 // Now print the reductions sorted by type. Each type will cause a clause 128 // like: reduction (+ : sum0, sum1, sum2) 129 for (const auto &ReductionClause : Clauses) { 130 str += " reduction ("; 131 str += MemoryAccess::getReductionOperatorStr(ReductionClause.first); 132 // Remove the first two symbols (", ") to make the output look pretty. 133 str += " : " + ReductionClause.second.substr(2) + ")"; 134 } 135 136 return str; 137 } 138 139 /// @brief Callback executed for each for node in the ast in order to print it. 140 static isl_printer *cbPrintFor(__isl_take isl_printer *Printer, 141 __isl_take isl_ast_print_options *Options, 142 __isl_keep isl_ast_node *Node, void *) { 143 144 const std::string BrokenReductionsStr = getBrokenReductionsStr(Node); 145 const std::string SimdPragmaStr = "#pragma simd"; 146 const std::string OmpPragmaStr = "#pragma omp parallel for"; 147 148 if (IslAstInfo::isInnermostParallel(Node)) 149 Printer = printLine(Printer, SimdPragmaStr + BrokenReductionsStr); 150 151 if (IslAstInfo::isOutermostParallel(Node)) 152 Printer = printLine(Printer, OmpPragmaStr + BrokenReductionsStr); 153 154 return isl_ast_node_for_print(Node, Printer, Options); 155 } 156 157 /// @brief Check if the current scheduling dimension is parallel 158 /// 159 /// In case the dimension is parallel we also check if any reduction 160 /// dependences is broken when we exploit this parallelism. If so, 161 /// @p IsReductionParallel will be set to true. The reduction dependences we use 162 /// to check are actually the union of the transitive closure of the initial 163 /// reduction dependences together with their reveresal. Even though these 164 /// dependences connect all iterations with each other (thus they are cyclic) 165 /// we can perform the parallelism check as we are only interested in a zero 166 /// (or non-zero) dependence distance on the dimension in question. 167 static bool astScheduleDimIsParallel(__isl_keep isl_ast_build *Build, 168 Dependences *D, 169 IslAstUserPayload *NodeInfo) { 170 if (!D->hasValidDependences()) 171 return false; 172 173 isl_union_map *Schedule = isl_ast_build_get_schedule(Build); 174 isl_union_map *Deps = D->getDependences( 175 Dependences::TYPE_RAW | Dependences::TYPE_WAW | Dependences::TYPE_WAR); 176 if (!D->isParallel(Schedule, Deps) && !isl_union_map_free(Schedule)) 177 return false; 178 179 isl_union_map *RedDeps = D->getDependences(Dependences::TYPE_TC_RED); 180 if (!D->isParallel(Schedule, RedDeps)) 181 NodeInfo->IsReductionParallel = true; 182 183 if (!NodeInfo->IsReductionParallel && !isl_union_map_free(Schedule)) 184 return true; 185 186 // Annotate reduction parallel nodes with the memory accesses which caused the 187 // reduction dependences parallel execution of the node conflicts with. 188 for (const auto &MaRedPair : D->getReductionDependences()) { 189 if (!MaRedPair.second) 190 continue; 191 RedDeps = isl_union_map_from_map(isl_map_copy(MaRedPair.second)); 192 if (!D->isParallel(Schedule, RedDeps)) 193 NodeInfo->BrokenReductions.insert(MaRedPair.first); 194 } 195 196 isl_union_map_free(Schedule); 197 return true; 198 } 199 200 // This method is executed before the construction of a for node. It creates 201 // an isl_id that is used to annotate the subsequently generated ast for nodes. 202 // 203 // In this function we also run the following analyses: 204 // 205 // - Detection of openmp parallel loops 206 // 207 static __isl_give isl_id *astBuildBeforeFor(__isl_keep isl_ast_build *Build, 208 void *User) { 209 AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User; 210 IslAstUserPayload *Payload = new IslAstUserPayload(); 211 isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", Payload); 212 Id = isl_id_set_free_user(Id, freeIslAstUserPayload); 213 BuildInfo->LastForNodeId = Id; 214 215 // Test for parallelism only if we are not already inside a parallel loop 216 if (!BuildInfo->InParallelFor) 217 BuildInfo->InParallelFor = Payload->IsOutermostParallel = 218 astScheduleDimIsParallel(Build, BuildInfo->Deps, Payload); 219 220 return Id; 221 } 222 223 // This method is executed after the construction of a for node. 224 // 225 // It performs the following actions: 226 // 227 // - Reset the 'InParallelFor' flag, as soon as we leave a for node, 228 // that is marked as openmp parallel. 229 // 230 static __isl_give isl_ast_node * 231 astBuildAfterFor(__isl_take isl_ast_node *Node, __isl_keep isl_ast_build *Build, 232 void *User) { 233 isl_id *Id = isl_ast_node_get_annotation(Node); 234 assert(Id && "Post order visit assumes annotated for nodes"); 235 IslAstUserPayload *Payload = (IslAstUserPayload *)isl_id_get_user(Id); 236 assert(Payload && "Post order visit assumes annotated for nodes"); 237 238 AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User; 239 assert(!Payload->Build && "Build environment already set"); 240 Payload->Build = isl_ast_build_copy(Build); 241 Payload->IsInnermost = (Id == BuildInfo->LastForNodeId); 242 243 // Innermost loops that are surrounded by parallel loops have not yet been 244 // tested for parallelism. Test them here to ensure we check all innermost 245 // loops for parallelism. 246 if (Payload->IsInnermost && BuildInfo->InParallelFor) { 247 if (Payload->IsOutermostParallel) 248 Payload->IsInnermostParallel = true; 249 else 250 Payload->IsInnermostParallel = 251 astScheduleDimIsParallel(Build, BuildInfo->Deps, Payload); 252 } else if (Payload->IsOutermostParallel) 253 BuildInfo->InParallelFor = false; 254 255 isl_id_free(Id); 256 return Node; 257 } 258 259 static __isl_give isl_ast_node *AtEachDomain(__isl_take isl_ast_node *Node, 260 __isl_keep isl_ast_build *Build, 261 void *User) { 262 assert(!isl_ast_node_get_annotation(Node) && "Node already annotated"); 263 264 IslAstUserPayload *Payload = new IslAstUserPayload(); 265 isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", Payload); 266 Id = isl_id_set_free_user(Id, freeIslAstUserPayload); 267 268 Payload->Build = isl_ast_build_copy(Build); 269 270 return isl_ast_node_set_annotation(Node, Id); 271 } 272 273 void IslAst::buildRunCondition(__isl_keep isl_ast_build *Build) { 274 // The conditions that need to be checked at run-time for this scop are 275 // available as an isl_set in the AssumedContext. We generate code for this 276 // check as follows. First, we generate an isl_pw_aff that is 1, if a certain 277 // combination of parameter values fulfills the conditions in the assumed 278 // context, and that is 0 otherwise. We then translate this isl_pw_aff into 279 // an isl_ast_expr. At run-time this expression can be evaluated and the 280 // optimized scop can be executed conditionally according to the result of the 281 // run-time check. 282 283 isl_aff *Zero = 284 isl_aff_zero_on_domain(isl_local_space_from_space(S->getParamSpace())); 285 isl_aff *One = 286 isl_aff_zero_on_domain(isl_local_space_from_space(S->getParamSpace())); 287 288 One = isl_aff_add_constant_si(One, 1); 289 290 isl_pw_aff *PwZero = isl_pw_aff_from_aff(Zero); 291 isl_pw_aff *PwOne = isl_pw_aff_from_aff(One); 292 293 PwOne = isl_pw_aff_intersect_domain(PwOne, S->getAssumedContext()); 294 PwZero = isl_pw_aff_intersect_domain( 295 PwZero, isl_set_complement(S->getAssumedContext())); 296 297 isl_pw_aff *Cond = isl_pw_aff_union_max(PwOne, PwZero); 298 299 RunCondition = isl_ast_build_expr_from_pw_aff(Build, Cond); 300 } 301 302 IslAst::IslAst(Scop *Scop, Dependences &D) : S(Scop) { 303 isl_ctx *Ctx = S->getIslCtx(); 304 isl_options_set_ast_build_atomic_upper_bound(Ctx, true); 305 isl_ast_build *Build; 306 AstBuildUserInfo BuildInfo; 307 308 if (UseContext) 309 Build = isl_ast_build_from_context(S->getContext()); 310 else 311 Build = isl_ast_build_from_context(isl_set_universe(S->getParamSpace())); 312 313 Build = isl_ast_build_set_at_each_domain(Build, AtEachDomain, nullptr); 314 315 isl_union_map *Schedule = 316 isl_union_map_intersect_domain(S->getSchedule(), S->getDomains()); 317 318 if (DetectParallel || PollyVectorizerChoice != VECTORIZER_NONE) { 319 BuildInfo.Deps = &D; 320 BuildInfo.InParallelFor = 0; 321 322 Build = isl_ast_build_set_before_each_for(Build, &astBuildBeforeFor, 323 &BuildInfo); 324 Build = 325 isl_ast_build_set_after_each_for(Build, &astBuildAfterFor, &BuildInfo); 326 } 327 328 buildRunCondition(Build); 329 330 Root = isl_ast_build_ast_from_schedule(Build, Schedule); 331 332 isl_ast_build_free(Build); 333 } 334 335 IslAst::~IslAst() { 336 isl_ast_node_free(Root); 337 isl_ast_expr_free(RunCondition); 338 } 339 340 __isl_give isl_ast_node *IslAst::getAst() { return isl_ast_node_copy(Root); } 341 __isl_give isl_ast_expr *IslAst::getRunCondition() { 342 return isl_ast_expr_copy(RunCondition); 343 } 344 345 void IslAstInfo::releaseMemory() { 346 if (Ast) { 347 delete Ast; 348 Ast = 0; 349 } 350 } 351 352 bool IslAstInfo::runOnScop(Scop &Scop) { 353 if (Ast) 354 delete Ast; 355 356 S = &Scop; 357 358 Dependences &D = getAnalysis<Dependences>(); 359 360 Ast = new IslAst(&Scop, D); 361 362 DEBUG(printScop(dbgs())); 363 return false; 364 } 365 366 __isl_give isl_ast_node *IslAstInfo::getAst() const { return Ast->getAst(); } 367 __isl_give isl_ast_expr *IslAstInfo::getRunCondition() const { 368 return Ast->getRunCondition(); 369 } 370 371 IslAstUserPayload *IslAstInfo::getNodePayload(__isl_keep isl_ast_node *Node) { 372 isl_id *Id = isl_ast_node_get_annotation(Node); 373 if (!Id) 374 return nullptr; 375 IslAstUserPayload *Payload = (IslAstUserPayload *)isl_id_get_user(Id); 376 isl_id_free(Id); 377 return Payload; 378 } 379 380 bool IslAstInfo::isInnermost(__isl_keep isl_ast_node *Node) { 381 IslAstUserPayload *Payload = getNodePayload(Node); 382 return Payload && Payload->IsInnermost; 383 } 384 385 bool IslAstInfo::isParallel(__isl_keep isl_ast_node *Node) { 386 return IslAstInfo::isInnermostParallel(Node) || 387 IslAstInfo::isOutermostParallel(Node); 388 } 389 390 bool IslAstInfo::isInnermostParallel(__isl_keep isl_ast_node *Node) { 391 IslAstUserPayload *Payload = getNodePayload(Node); 392 return Payload && Payload->IsInnermostParallel; 393 } 394 395 bool IslAstInfo::isOutermostParallel(__isl_keep isl_ast_node *Node) { 396 IslAstUserPayload *Payload = getNodePayload(Node); 397 return Payload && Payload->IsOutermostParallel; 398 } 399 400 bool IslAstInfo::isReductionParallel(__isl_keep isl_ast_node *Node) { 401 IslAstUserPayload *Payload = getNodePayload(Node); 402 return Payload && Payload->IsReductionParallel; 403 } 404 405 isl_union_map *IslAstInfo::getSchedule(__isl_keep isl_ast_node *Node) { 406 IslAstUserPayload *Payload = getNodePayload(Node); 407 return Payload ? isl_ast_build_get_schedule(Payload->Build) : nullptr; 408 } 409 410 IslAstInfo::MemoryAccessSet * 411 IslAstInfo::getBrokenReductions(__isl_keep isl_ast_node *Node) { 412 IslAstUserPayload *Payload = getNodePayload(Node); 413 return Payload ? &Payload->BrokenReductions : nullptr; 414 } 415 416 isl_ast_build *IslAstInfo::getBuild(__isl_keep isl_ast_node *Node) { 417 IslAstUserPayload *Payload = getNodePayload(Node); 418 return Payload ? Payload->Build : nullptr; 419 } 420 421 void IslAstInfo::printScop(raw_ostream &OS) const { 422 isl_ast_print_options *Options; 423 isl_ast_node *RootNode = getAst(); 424 isl_ast_expr *RunCondition = getRunCondition(); 425 char *RtCStr, *AstStr; 426 427 Scop &S = getCurScop(); 428 Options = isl_ast_print_options_alloc(S.getIslCtx()); 429 Options = isl_ast_print_options_set_print_for(Options, cbPrintFor, nullptr); 430 431 isl_printer *P = isl_printer_to_str(S.getIslCtx()); 432 P = isl_printer_print_ast_expr(P, RunCondition); 433 RtCStr = isl_printer_get_str(P); 434 P = isl_printer_flush(P); 435 P = isl_printer_indent(P, 4); 436 P = isl_printer_set_output_format(P, ISL_FORMAT_C); 437 P = isl_ast_node_print(RootNode, P, Options); 438 AstStr = isl_printer_get_str(P); 439 440 Function *F = S.getRegion().getEntry()->getParent(); 441 isl_union_map *Schedule = 442 isl_union_map_intersect_domain(S.getSchedule(), S.getDomains()); 443 444 OS << ":: isl ast :: " << F->getName() << " :: " << S.getRegion().getNameStr() 445 << "\n"; 446 DEBUG(dbgs() << S.getContextStr() << "\n"; isl_union_map_dump(Schedule)); 447 OS << "\nif (" << RtCStr << ")\n\n"; 448 OS << AstStr << "\n"; 449 OS << "else\n"; 450 OS << " { /* original code */ }\n\n"; 451 452 isl_ast_expr_free(RunCondition); 453 isl_union_map_free(Schedule); 454 isl_ast_node_free(RootNode); 455 isl_printer_free(P); 456 } 457 458 void IslAstInfo::getAnalysisUsage(AnalysisUsage &AU) const { 459 // Get the Common analysis usage of ScopPasses. 460 ScopPass::getAnalysisUsage(AU); 461 AU.addRequired<ScopInfo>(); 462 AU.addRequired<Dependences>(); 463 } 464 465 char IslAstInfo::ID = 0; 466 467 Pass *polly::createIslAstInfoPass() { return new IslAstInfo(); } 468 469 INITIALIZE_PASS_BEGIN(IslAstInfo, "polly-ast", 470 "Polly - Generate an AST of the SCoP (isl)", false, 471 false); 472 INITIALIZE_PASS_DEPENDENCY(ScopInfo); 473 INITIALIZE_PASS_DEPENDENCY(Dependences); 474 INITIALIZE_PASS_END(IslAstInfo, "polly-ast", 475 "Polly - Generate an AST from the SCoP (isl)", false, false) 476