IslAst.cpp 29 KB

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