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