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 using namespace llvm; 38 using namespace polly; 39 40 #define DEBUG_TYPE "polly-ast" 41 42 static cl::opt<bool> UseContext("polly-ast-use-context", 43 cl::desc("Use context"), cl::Hidden, 44 cl::init(false), cl::ZeroOrMore, 45 cl::cat(PollyCategory)); 46 47 static cl::opt<bool> DetectParallel("polly-ast-detect-parallel", 48 cl::desc("Detect parallelism"), cl::Hidden, 49 cl::init(false), cl::ZeroOrMore, 50 cl::cat(PollyCategory)); 51 52 namespace polly { 53 class IslAst { 54 public: 55 IslAst(Scop *Scop, Dependences &D); 56 57 ~IslAst(); 58 59 /// Print a source code representation of the program. 60 void pprint(llvm::raw_ostream &OS); 61 62 __isl_give isl_ast_node *getAst(); 63 64 /// @brief Get the run-time conditions for the Scop. 65 __isl_give isl_ast_expr *getRunCondition(); 66 67 private: 68 Scop *S; 69 isl_ast_node *Root; 70 isl_ast_expr *RunCondition; 71 72 void buildRunCondition(__isl_keep isl_ast_build *Context); 73 }; 74 } // End namespace polly. 75 76 // Temporary information used when building the ast. 77 struct AstBuildUserInfo { 78 // The dependence information. 79 Dependences *Deps; 80 81 // We are inside a parallel for node. 82 int InParallelFor; 83 }; 84 85 // Print a loop annotated with OpenMP or vector pragmas. 86 static __isl_give isl_printer * 87 printParallelFor(__isl_keep isl_ast_node *Node, __isl_take isl_printer *Printer, 88 __isl_take isl_ast_print_options *PrintOptions, 89 IslAstUserPayload *Info) { 90 if (Info) { 91 if (Info->IsInnermostParallel) { 92 Printer = isl_printer_start_line(Printer); 93 Printer = isl_printer_print_str(Printer, "#pragma simd"); 94 if (Info->IsReductionParallel) 95 Printer = isl_printer_print_str(Printer, " reduction"); 96 Printer = isl_printer_end_line(Printer); 97 } 98 if (Info->IsOutermostParallel) { 99 Printer = isl_printer_start_line(Printer); 100 Printer = isl_printer_print_str(Printer, "#pragma omp parallel for"); 101 if (Info->IsReductionParallel) 102 Printer = isl_printer_print_str(Printer, " reduction"); 103 Printer = isl_printer_end_line(Printer); 104 } 105 } 106 return isl_ast_node_for_print(Node, Printer, PrintOptions); 107 } 108 109 // Print an isl_ast_for. 110 static __isl_give isl_printer * 111 printFor(__isl_take isl_printer *Printer, 112 __isl_take isl_ast_print_options *PrintOptions, 113 __isl_keep isl_ast_node *Node, void *User) { 114 isl_id *Id = isl_ast_node_get_annotation(Node); 115 if (!Id) 116 return isl_ast_node_for_print(Node, Printer, PrintOptions); 117 118 struct IslAstUserPayload *Info = 119 (struct IslAstUserPayload *)isl_id_get_user(Id); 120 Printer = printParallelFor(Node, Printer, PrintOptions, Info); 121 isl_id_free(Id); 122 return Printer; 123 } 124 125 // Allocate an AstNodeInfo structure and initialize it with default values. 126 static struct IslAstUserPayload *allocateIslAstUser() { 127 struct IslAstUserPayload *NodeInfo; 128 NodeInfo = 129 (struct IslAstUserPayload *)malloc(sizeof(struct IslAstUserPayload)); 130 NodeInfo->Context = 0; 131 NodeInfo->IsOutermostParallel = 0; 132 NodeInfo->IsInnermostParallel = 0; 133 NodeInfo->IsReductionParallel = false; 134 return NodeInfo; 135 } 136 137 // Free the AstNodeInfo structure. 138 static void freeIslAstUser(void *Ptr) { 139 struct IslAstUserPayload *UserStruct = (struct IslAstUserPayload *)Ptr; 140 isl_ast_build_free(UserStruct->Context); 141 free(UserStruct); 142 } 143 144 // Check if the current scheduling dimension is parallel. 145 // 146 // We check for parallelism by verifying that the loop does not carry any 147 // dependences. 148 // 149 // Parallelism test: if the distance is zero in all outer dimensions, then it 150 // has to be zero in the current dimension as well. 151 // 152 // Implementation: first, translate dependences into time space, then force 153 // outer dimensions to be equal. If the distance is zero in the current 154 // dimension, then the loop is parallel. The distance is zero in the current 155 // dimension if it is a subset of a map with equal values for the current 156 // dimension. 157 static bool astScheduleDimIsParallel(__isl_keep isl_ast_build *Build, 158 __isl_take isl_union_map *Deps) { 159 isl_union_map *Schedule; 160 isl_map *ScheduleDeps, *Test; 161 isl_space *ScheduleSpace; 162 unsigned Dimension, IsParallel; 163 164 Schedule = isl_ast_build_get_schedule(Build); 165 ScheduleSpace = isl_ast_build_get_schedule_space(Build); 166 167 Dimension = isl_space_dim(ScheduleSpace, isl_dim_out) - 1; 168 169 Deps = isl_union_map_apply_range(Deps, isl_union_map_copy(Schedule)); 170 Deps = isl_union_map_apply_domain(Deps, Schedule); 171 172 if (isl_union_map_is_empty(Deps)) { 173 isl_union_map_free(Deps); 174 isl_space_free(ScheduleSpace); 175 return true; 176 } 177 178 ScheduleDeps = isl_map_from_union_map(Deps); 179 180 for (unsigned i = 0; i < Dimension; i++) 181 ScheduleDeps = isl_map_equate(ScheduleDeps, isl_dim_out, i, isl_dim_in, i); 182 183 Test = isl_map_universe(isl_map_get_space(ScheduleDeps)); 184 Test = isl_map_equate(Test, isl_dim_out, Dimension, isl_dim_in, Dimension); 185 IsParallel = isl_map_is_subset(ScheduleDeps, Test); 186 187 isl_space_free(ScheduleSpace); 188 isl_map_free(Test); 189 isl_map_free(ScheduleDeps); 190 191 return IsParallel; 192 } 193 194 /// @brief Check if the current scheduling dimension is parallel 195 /// 196 /// In case the dimension is parallel we also check if any reduction 197 /// dependences is broken when we exploit this parallelism. If so, 198 /// @p IsReductionParallel will be set to true. The reduction dependences we use 199 /// to check are actually the union of the transitive closure of the initial 200 /// reduction dependences together with their reveresal. Even though these 201 /// 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 zero 203 /// (or non-zero) dependence distance on the dimension in question. 204 static bool astScheduleDimIsParallel(__isl_keep isl_ast_build *Build, 205 Dependences *D, 206 bool &IsReductionParallel) { 207 if (!D->hasValidDependences()) 208 return false; 209 210 isl_union_map *Deps = D->getDependences( 211 Dependences::TYPE_RAW | Dependences::TYPE_WAW | Dependences::TYPE_WAR); 212 if (!astScheduleDimIsParallel(Build, Deps)) 213 return false; 214 215 isl_union_map *RedDeps = D->getDependences(Dependences::TYPE_TC_RED); 216 if (!astScheduleDimIsParallel(Build, RedDeps)) 217 IsReductionParallel = true; 218 219 return true; 220 } 221 222 // Mark a for node openmp parallel, if it is the outermost parallel for node. 223 static void markOpenmpParallel(__isl_keep isl_ast_build *Build, 224 struct AstBuildUserInfo *BuildInfo, 225 struct IslAstUserPayload *NodeInfo) { 226 if (BuildInfo->InParallelFor) 227 return; 228 229 if (astScheduleDimIsParallel(Build, BuildInfo->Deps, 230 NodeInfo->IsReductionParallel)) { 231 BuildInfo->InParallelFor = 1; 232 NodeInfo->IsOutermostParallel = 1; 233 } 234 } 235 236 // This method is executed before the construction of a for node. It creates 237 // an isl_id that is used to annotate the subsequently generated ast for nodes. 238 // 239 // In this function we also run the following analyses: 240 // 241 // - Detection of openmp parallel loops 242 // 243 static __isl_give isl_id *astBuildBeforeFor(__isl_keep isl_ast_build *Build, 244 void *User) { 245 struct AstBuildUserInfo *BuildInfo = (struct AstBuildUserInfo *)User; 246 struct IslAstUserPayload *NodeInfo = allocateIslAstUser(); 247 isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", NodeInfo); 248 Id = isl_id_set_free_user(Id, freeIslAstUser); 249 250 markOpenmpParallel(Build, BuildInfo, NodeInfo); 251 252 return Id; 253 } 254 255 // Returns 0 when Node contains loops, otherwise returns -1. This search 256 // function uses ISL's way to iterate over lists of isl_ast_nodes with 257 // isl_ast_node_list_foreach. Please use the single argument wrapper function 258 // that returns a bool instead of using this function directly. 259 static int containsLoops(__isl_take isl_ast_node *Node, void *User) { 260 if (!Node) 261 return -1; 262 263 switch (isl_ast_node_get_type(Node)) { 264 case isl_ast_node_for: 265 isl_ast_node_free(Node); 266 return 0; 267 case isl_ast_node_block: { 268 isl_ast_node_list *List = isl_ast_node_block_get_children(Node); 269 int Res = isl_ast_node_list_foreach(List, &containsLoops, nullptr); 270 isl_ast_node_list_free(List); 271 isl_ast_node_free(Node); 272 return Res; 273 } 274 case isl_ast_node_if: { 275 int Res = -1; 276 if (0 == containsLoops(isl_ast_node_if_get_then(Node), nullptr) || 277 (isl_ast_node_if_has_else(Node) && 278 0 == containsLoops(isl_ast_node_if_get_else(Node), nullptr))) 279 Res = 0; 280 isl_ast_node_free(Node); 281 return Res; 282 } 283 case isl_ast_node_user: 284 default: 285 isl_ast_node_free(Node); 286 return -1; 287 } 288 } 289 290 // Returns true when Node contains loops. 291 static bool containsLoops(__isl_take isl_ast_node *Node) { 292 return 0 == containsLoops(Node, nullptr); 293 } 294 295 // This method is executed after the construction of a for node. 296 // 297 // It performs the following actions: 298 // 299 // - Reset the 'InParallelFor' flag, as soon as we leave a for node, 300 // that is marked as openmp parallel. 301 // 302 static __isl_give isl_ast_node * 303 astBuildAfterFor(__isl_take isl_ast_node *Node, __isl_keep isl_ast_build *Build, 304 void *User) { 305 isl_id *Id = isl_ast_node_get_annotation(Node); 306 if (!Id) 307 return Node; 308 struct IslAstUserPayload *Info = 309 (struct IslAstUserPayload *)isl_id_get_user(Id); 310 struct AstBuildUserInfo *BuildInfo = (struct AstBuildUserInfo *)User; 311 312 if (Info) { 313 if (Info->IsOutermostParallel) 314 BuildInfo->InParallelFor = 0; 315 if (!containsLoops(isl_ast_node_for_get_body(Node))) 316 if (astScheduleDimIsParallel(Build, BuildInfo->Deps, 317 Info->IsReductionParallel)) 318 Info->IsInnermostParallel = 1; 319 if (!Info->Context) 320 Info->Context = isl_ast_build_copy(Build); 321 } 322 323 isl_id_free(Id); 324 return Node; 325 } 326 327 static __isl_give isl_ast_node *AtEachDomain(__isl_take isl_ast_node *Node, 328 __isl_keep isl_ast_build *Context, 329 void *User) { 330 struct IslAstUserPayload *Info = nullptr; 331 isl_id *Id = isl_ast_node_get_annotation(Node); 332 333 if (Id) 334 Info = (struct IslAstUserPayload *)isl_id_get_user(Id); 335 336 if (!Info) { 337 // Allocate annotations once: parallel for detection might have already 338 // allocated the annotations for this node. 339 Info = allocateIslAstUser(); 340 Id = isl_id_alloc(isl_ast_node_get_ctx(Node), nullptr, Info); 341 Id = isl_id_set_free_user(Id, &freeIslAstUser); 342 } 343 344 if (!Info->Context) 345 Info->Context = isl_ast_build_copy(Context); 346 347 return isl_ast_node_set_annotation(Node, Id); 348 } 349 350 void IslAst::buildRunCondition(__isl_keep isl_ast_build *Context) { 351 // The conditions that need to be checked at run-time for this scop are 352 // available as an isl_set in the AssumedContext. We generate code for this 353 // check as follows. First, we generate an isl_pw_aff that is 1, if a certain 354 // combination of parameter values fulfills the conditions in the assumed 355 // context, and that is 0 otherwise. We then translate this isl_pw_aff into 356 // an isl_ast_expr. At run-time this expression can be evaluated and the 357 // optimized scop can be executed conditionally according to the result of the 358 // run-time check. 359 360 isl_aff *Zero = 361 isl_aff_zero_on_domain(isl_local_space_from_space(S->getParamSpace())); 362 isl_aff *One = 363 isl_aff_zero_on_domain(isl_local_space_from_space(S->getParamSpace())); 364 365 One = isl_aff_add_constant_si(One, 1); 366 367 isl_pw_aff *PwZero = isl_pw_aff_from_aff(Zero); 368 isl_pw_aff *PwOne = isl_pw_aff_from_aff(One); 369 370 PwOne = isl_pw_aff_intersect_domain(PwOne, S->getAssumedContext()); 371 PwZero = isl_pw_aff_intersect_domain( 372 PwZero, isl_set_complement(S->getAssumedContext())); 373 374 isl_pw_aff *Cond = isl_pw_aff_union_max(PwOne, PwZero); 375 376 RunCondition = isl_ast_build_expr_from_pw_aff(Context, Cond); 377 } 378 379 IslAst::IslAst(Scop *Scop, Dependences &D) : S(Scop) { 380 isl_ctx *Ctx = S->getIslCtx(); 381 isl_options_set_ast_build_atomic_upper_bound(Ctx, true); 382 isl_ast_build *Context; 383 struct AstBuildUserInfo BuildInfo; 384 385 if (UseContext) 386 Context = isl_ast_build_from_context(S->getContext()); 387 else 388 Context = isl_ast_build_from_context(isl_set_universe(S->getParamSpace())); 389 390 Context = isl_ast_build_set_at_each_domain(Context, AtEachDomain, nullptr); 391 392 isl_union_map *Schedule = 393 isl_union_map_intersect_domain(S->getSchedule(), S->getDomains()); 394 395 Function *F = Scop->getRegion().getEntry()->getParent(); 396 (void)F; 397 398 DEBUG(dbgs() << ":: isl ast :: " << F->getName() 399 << " :: " << Scop->getRegion().getNameStr() << "\n"); 400 401 DEBUG(dbgs() << S->getContextStr() << "\n"; isl_union_map_dump(Schedule)); 402 403 if (DetectParallel || PollyVectorizerChoice != VECTORIZER_NONE) { 404 BuildInfo.Deps = &D; 405 BuildInfo.InParallelFor = 0; 406 407 Context = isl_ast_build_set_before_each_for(Context, &astBuildBeforeFor, 408 &BuildInfo); 409 Context = isl_ast_build_set_after_each_for(Context, &astBuildAfterFor, 410 &BuildInfo); 411 } 412 413 buildRunCondition(Context); 414 415 Root = isl_ast_build_ast_from_schedule(Context, Schedule); 416 417 isl_ast_build_free(Context); 418 419 DEBUG(pprint(dbgs())); 420 } 421 422 IslAst::~IslAst() { 423 isl_ast_node_free(Root); 424 isl_ast_expr_free(RunCondition); 425 } 426 427 /// Print a C like representation of the program. 428 void IslAst::pprint(llvm::raw_ostream &OS) { 429 isl_ast_node *Root; 430 isl_ast_print_options *Options; 431 432 Options = isl_ast_print_options_alloc(S->getIslCtx()); 433 Options = isl_ast_print_options_set_print_for(Options, &printFor, nullptr); 434 435 isl_printer *P = isl_printer_to_str(S->getIslCtx()); 436 P = isl_printer_set_output_format(P, ISL_FORMAT_C); 437 438 P = isl_printer_print_ast_expr(P, RunCondition); 439 char *result = isl_printer_get_str(P); 440 P = isl_printer_flush(P); 441 442 OS << "\nif (" << result << ")\n\n"; 443 P = isl_printer_indent(P, 4); 444 445 Root = getAst(); 446 P = isl_ast_node_print(Root, P, Options); 447 result = isl_printer_get_str(P); 448 OS << result << "\n"; 449 OS << "else\n"; 450 OS << " { /* original code */ }\n\n"; 451 isl_printer_free(P); 452 isl_ast_node_free(Root); 453 } 454 455 __isl_give isl_ast_node *IslAst::getAst() { return isl_ast_node_copy(Root); } 456 __isl_give isl_ast_expr *IslAst::getRunCondition() { 457 return isl_ast_expr_copy(RunCondition); 458 } 459 460 void IslAstInfo::pprint(llvm::raw_ostream &OS) { Ast->pprint(OS); } 461 462 void IslAstInfo::releaseMemory() { 463 if (Ast) { 464 delete Ast; 465 Ast = 0; 466 } 467 } 468 469 bool IslAstInfo::runOnScop(Scop &Scop) { 470 if (Ast) 471 delete Ast; 472 473 S = &Scop; 474 475 Dependences &D = getAnalysis<Dependences>(); 476 477 Ast = new IslAst(&Scop, D); 478 479 return false; 480 } 481 482 __isl_give isl_ast_node *IslAstInfo::getAst() { return Ast->getAst(); } 483 __isl_give isl_ast_expr *IslAstInfo::getRunCondition() { 484 return Ast->getRunCondition(); 485 } 486 487 void IslAstInfo::printScop(raw_ostream &OS) const { 488 Function *F = S->getRegion().getEntry()->getParent(); 489 490 OS << F->getName() << "():\n"; 491 492 Ast->pprint(OS); 493 } 494 495 IslAstUserPayload *IslAstInfo::getNodePayload(__isl_keep isl_ast_node *Node) { 496 isl_id *Id = isl_ast_node_get_annotation(Node); 497 if (!Id) 498 return nullptr; 499 IslAstUserPayload *Payload = (IslAstUserPayload *)isl_id_get_user(Id); 500 isl_id_free(Id); 501 return Payload; 502 } 503 504 bool IslAstInfo::isParallel(__isl_keep isl_ast_node *Node) { 505 return (isInnermostParallel(Node) || isOuterParallel(Node)) && 506 !isReductionParallel(Node); 507 } 508 509 bool IslAstInfo::isInnermostParallel(__isl_keep isl_ast_node *Node) { 510 IslAstUserPayload *Payload = getNodePayload(Node); 511 return Payload && Payload->IsInnermostParallel && 512 !Payload->IsReductionParallel; 513 } 514 515 bool IslAstInfo::isOuterParallel(__isl_keep isl_ast_node *Node) { 516 IslAstUserPayload *Payload = getNodePayload(Node); 517 return Payload && Payload->IsOutermostParallel && 518 !Payload->IsReductionParallel; 519 } 520 521 bool IslAstInfo::isReductionParallel(__isl_keep isl_ast_node *Node) { 522 IslAstUserPayload *Payload = getNodePayload(Node); 523 return Payload && Payload->IsReductionParallel; 524 } 525 526 void IslAstInfo::getAnalysisUsage(AnalysisUsage &AU) const { 527 // Get the Common analysis usage of ScopPasses. 528 ScopPass::getAnalysisUsage(AU); 529 AU.addRequired<ScopInfo>(); 530 AU.addRequired<Dependences>(); 531 } 532 533 char IslAstInfo::ID = 0; 534 535 Pass *polly::createIslAstInfoPass() { return new IslAstInfo(); } 536 537 INITIALIZE_PASS_BEGIN(IslAstInfo, "polly-ast", 538 "Polly - Generate an AST of the SCoP (isl)", false, 539 false); 540 INITIALIZE_PASS_DEPENDENCY(ScopInfo); 541 INITIALIZE_PASS_DEPENDENCY(Dependences); 542 INITIALIZE_PASS_END(IslAstInfo, "polly-ast", 543 "Polly - Generate an AST from the SCoP (isl)", false, false) 544