summaryrefslogtreecommitdiffstats
path: root/polly/lib/Support
Commit message (Collapse)AuthorAgeFilesLines
...
* Delinearize memory accesses that reference parameters coming from function callsTobias Grosser2017-05-271-0/+35
| | | | | | | | | | | | | | Certain affine memory accesses which we model today might contain products of parameters which we might combined into a new parameter to be able to create an affine expression that represents these memory accesses. Especially in the context of OpenCL, this approach looses information as memory accesses such as A[get_global_id(0) * N + get_global_id(1)] are assumed to be linear. We correctly recover their multi-dimensional structure by assuming that parameters that are the result of a function call at IR level likely are not parameters, but indeed induction variables. The resulting access is now A[get_global_id(0)][get_global_id(1)] for an array A[][N]. llvm-svn: 304075
* Allow side-effect free function calls in valid affine SCEVsTobias Grosser2017-05-271-0/+29
| | | | | | | | | | | | | | | Side-effect free function calls with only constant parameters can be easily re-generated and consequently do not prevent us from modeling a SCEV. This change allows array subscripts to reference function calls such as 'get_global_id()' as used in OpenCL. We use the function name plus the constant operands to name the parameter. This is possible as the function name is required and is not dropped in release builds the same way names of llvm::Values are dropped. We also provide more readable names for common OpenCL functions, to make it easy to understand the polyhedral model we generate. llvm-svn: 304074
* [Polly] Add handling of Top Level RegionsPhilip Pfaffe2017-05-241-3/+9
| | | | | | | | | | | | | | | | | | | | | | | | | | | Summary: My goal is to make the newly added `AllowWholeFunctions` options more usable/powerful. The changes to ScopBuilder.cpp are exclusively checks to prevent `Region.getExit()` from being dereferenced, since Top Level Regions (TLRs) don't have an exit block. In ScopDetection's `isValidCFG`, I removed a check that disallowed ReturnInstructions to have return values. This might of course have been intentional, so I would welcome your feedback on this and maybe a small explanation why return values are forbidden. Maybe it can be done but needs more changes elsewhere? The remaining changes in ScopDetection are simply to consider the AllowWholeFunctions option in more places, i.e. allow TLRs when it is set and once again avoid derefererncing `getExit()` if it doesn't exist. Finally, in ScopHelper.cpp I extended `polly::isErrorBlock` to handle regions without exit blocks as well: The original check was if a given BasicBlock dominates all predecessors of the exit block. Therefore I do the same for TLRs by regarding all BasicBlocks terminating with a ReturnInst as predecessors of a "virtual" function exit block. Patch by: Lukas Boehm Reviewers: philip.pfaffe, grosser, Meinersbur Reviewed By: grosser Subscribers: pollydev, llvm-commits, bollu Tags: #polly Differential Revision: https://reviews.llvm.org/D33411 llvm-svn: 303790
* [Polly][NewPM] Port IslAst to the new ScopPassManagerPhilip Pfaffe2017-05-231-2/+2
| | | | | | | | | | | | | | | | Summary: This patch ports IslAst to the new PM. The change is mostly straightforward. The only major modification required is making IslAst move-only, to correctly manage the isl resources it owns. Reviewers: grosser, Meinersbur Reviewed By: grosser Subscribers: nemanjai, pollydev, llvm-commits Tags: #polly Differential Revision: https://reviews.llvm.org/D33422 llvm-svn: 303622
* Adjust formattingTobias Grosser2017-05-151-1/+1
| | | | llvm-svn: 303065
* [Polly] Fix code generation of llvm.expect intrinsicTobias Grosser2017-05-141-1/+0
| | | | | | | | | | | | | | | | | | | At the time of code generation, an instruction with an llvm intrinsic is ignored in copyBB. However, if the value of the instruction is used later in the program, the value needs to be synthesized. However, this is causing some issues with the instructions being generated in a hoisted basic block. Removing llvm.expect from the list of ignored intrinsics fixes this bug. This resolves http://llvm.org/PR32324. Contributed-by: Annanay Agarwal <cs14btech11001@iith.ac.in> Tags: #polly Differential Revision: https://reviews.llvm.org/D32992 llvm-svn: 303006
* [Polly][NewPM] Port ScopDetection to the new PassManagerPhilip Pfaffe2017-05-121-2/+2
| | | | | | | | | | | | | | | | Summary: This is a proof of concept of how to port polly-passes to the new PassManager architecture. This approach works ootb for Function-Passes, but might not be directly applicable to Scop/Region-Passes. While we could just run the Analyses/Transforms over functions instead, we'd surrender the nice pipelining behaviour we have now. Reviewers: Meinersbur, grosser Reviewed By: grosser Subscribers: pollydev, sanjoy, nemanjai, llvm-commits Tags: #polly Differential Revision: https://reviews.llvm.org/D31459 llvm-svn: 302902
* [VirtualInstruction] Do a lookup instead of a linear search. NFC.Michael Kruse2017-05-111-20/+1
| | | | llvm-svn: 302837
* [Polly] Added OpenCL Runtime to GPURuntime Library for GPGPU CodeGenSiddharth Bhat2017-05-071-1/+20
| | | | | | | | | | | | | | | | | | | | | | | Summary: When compiling for GPU, one can now choose to compile for OpenCL or CUDA, with the corresponding polly-gpu-runtime flag (libopencl / libcudart). The GPURuntime library (GPUJIT) has been extended with the OpenCL Runtime library for that purpose, correctly choosing the corresponding library calls to the option chosen when compiling (via different initialization calls). Additionally, a specific GPU Target architecture can now be chosen with -polly-gpu-arch (only nvptx64 implemented thus far). Reviewers: grosser, bollu, Meinersbur, etherzhhb, singam-sanjay Reviewed By: grosser, Meinersbur Subscribers: singam-sanjay, llvm-commits, pollydev, nemanjai, mgorny, yaxunl, Anastasia Tags: #polly Differential Revision: https://reviews.llvm.org/D32431 llvm-svn: 302379
* [DeLICM] Known knowledge.Michael Kruse2017-05-061-0/+1
| | | | | | | | | | | | | Extend the Knowledge class to store information about the contents of array elements and which values are written. Two knowledges do not conflict the known content is the same. The content information if computed from writes to and loads from the array elements, and represented by "ValInst": isl spaces that compare equal if the value represented is the same. Differential Revision: https://reviews.llvm.org/D31247 llvm-svn: 302339
* Revert "[Polly] Added OpenCL Runtime to GPURuntime Library for GPGPU CodeGen"Siddharth Bhat2017-05-051-20/+1
| | | | | | | | | | | | | | | This reverts commit 17a84e414adb51ee375d14836d4c2a817b191933. Patches should have been submitted in the order of: 1. D32852 2. D32854 3. D32431 I mistakenly pushed D32431(3) first. Reverting to push in the correct order. llvm-svn: 302217
* [Polly] Added OpenCL Runtime to GPURuntime Library for GPGPU CodeGenSiddharth Bhat2017-05-051-1/+20
| | | | | | | | | | | | | | | | | | | | | | | Summary: When compiling for GPU, one can now choose to compile for OpenCL or CUDA, with the corresponding polly-gpu-runtime flag (libopencl / libcudart). The GPURuntime library (GPUJIT) has been extended with the OpenCL Runtime library for that purpose, correctly choosing the corresponding library calls to the option chosen when compiling (via different initialization calls). Additionally, a specific GPU Target architecture can now be chosen with -polly-gpu-arch (only nvptx64 implemented thus far). Reviewers: grosser, bollu, Meinersbur, etherzhhb, singam-sanjay Reviewed By: grosser, Meinersbur Subscribers: singam-sanjay, llvm-commits, pollydev, nemanjai, mgorny, yaxunl, Anastasia Tags: #polly Differential Revision: https://reviews.llvm.org/D32431 llvm-svn: 302215
* Introduce VirtualUse. NFC.Michael Kruse2017-05-041-0/+146
| | | | | | | | | | | | | | | | | | | If a ScopStmt references a (scalar) value, there are multiple possibilities where this value can come. The decision about what kind of use it is must be handled consistently at different places, which can be error-prone. VirtualUse is meant to centralize the handling of the different types of value uses. This patch makes ScopBuilder and CodeGeneration use VirtualUse. This already helps to show inconsistencies with the value handling. In order to keep this patch NFC, exceptions to the general rules are added. These might be fixed later if they turn to problems. Overall, this should result in fewer post-codegen IR-verification errors, but instead assertion failures in `getNewValue` that are closer to the actual error. Differential Revision: https://reviews.llvm.org/D32667 llvm-svn: 302157
* [ScopInfo] Do not use LLVM names to identify statements, arrays, and parametersTobias Grosser2017-05-031-7/+26
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | LLVM-IR names are commonly available in debug builds, but often not in release builds. Hence, using LLVM-IR names to identify statements or memory reference results makes the behavior of Polly depend on the compile mode. This is undesirable. Hence, we now just number the statements instead of using LLVM-IR names to identify them (this issue has previously been brought up by Zino Benaissa). However, as LLVM-IR names help in making test cases more readable, we add an option '-polly-use-llvm-names' to still use LLVM-IR names. This flag is by default set in the polly tests to make test cases more readable. This change reduces the time in ScopInfo from 32 seconds to 2 seconds for the following test case provided by Eli Friedman <efriedma@codeaurora.org> (already used in one of the previous commits): struct X { int x; }; void a(); #define SIG (int x, X **y, X **z) typedef void (*fn)SIG; #define FN { for (int i = 0; i < x; ++i) { (*y)[i].x += (*z)[i].x; } a(); } #define FN5 FN FN FN FN FN #define FN25 FN5 FN5 FN5 FN5 #define FN125 FN25 FN25 FN25 FN25 FN25 #define FN250 FN125 FN125 #define FN1250 FN250 FN250 FN250 FN250 FN250 void x SIG { FN1250 } For a larger benchmark I have on-hand (10000 loops), this reduces the time for running -polly-scops from 5 minutes to 4 minutes, a reduction by 20%. The reason for this large speedup is that our previous use of printAsOperand had a quadratic cost, as for each printed and unnamed operand the full function was scanned to find the instruction number that identifies the operand. We do not need to adjust the way memory reference ids are constructured, as they do not use LLVM values. Reviewed by: efriedma Tags: #polly Differential Revision: https://reviews.llvm.org/D32789 llvm-svn: 302072
* Use isl C++ foreach implementationTobias Grosser2017-04-142-111/+16
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | This commit switches Polly over to the isl::obj::foreach_* implementation, which is part of the new isl bindings and follows the foreach pattern established in Polly by Michael Kruse. The original isl C function: isl_stat isl_union_set_foreach_set(__isl_keep isl_union_set *uset, isl_stat (*fn)(__isl_take isl_set *set, void *user), void *user); which required the user to define a static callback function to which all interesting parameters are passed via a 'void *' user-pointer, is on the C++ side available as a function that takes a std::function<>, which can carry any additional arguments without the need for a user pointer: stat UnionSet::foreach_set(const std::function<stat(set)> &fn) const; The following code illustrates the use of the new C++ interface: auto Lambda = [=, &Result](isl::set Set) -> isl::stat { auto Shifted = shiftDimension(Set, Pos, Amount); Result = Result.add(Shifted); return isl::stat::ok; } UnionSet.foreach_set(Lambda); Polly had some specialized foreach functions which did not require the lambdas to return a status flag. We remove these functions in this commit to move Polly completely over to the new isl interface. We may in the future discuss if functors without return values can be supported easily. Another extension proposed by Michael Kruse is the use of C++ iterators to allow the use of normal for loops to iterate over these sets. Such an extension would allow us to further simplify the code. Reviewed-by: Michael Kruse <llvm@meinersbur.de> Differential Revision: https://reviews.llvm.org/D30620 llvm-svn: 300323
* [Support] Add functions to ISLTools.Michael Kruse2017-03-221-0/+139
| | | | | | | | | | | Add shiftDim and convertZoneToTimepoints overloads for isl maps. Add distributeDomain, liftDomains and applyDomainRange functions. These are going to be used in https://reviews.llvm.org/D31247 (Add known array contents to Knowledge) llvm-svn: 298543
* [ScopInfo/PruneUnprofitable] Move default profitability check.Michael Kruse2017-03-171-1/+1
| | | | | | | | | | | | | | | | | | | | In the previous default ScopInfo applied the profitability heuristic for scalar accesses (-polly-unprofitable-scalar-accs=true) and the -polly-prune-unprofitable was disabled by default (-polly-enable-prune-unprofitable=false) as that pruning was already done. This changes switches the defaults to -polly-unprofitable-scalar-accs=true -polly-enable-prune-unprofitable=false such that the scalar access heuristic check is done by the pass. This allows passes between ScopInfo and PruneUnprofitable to optimize away scalar accesses. Without enabling such intermediate passes, there is no change in behaviour of profitability checks in a PassManagerBuilder built pass chain, but it allows us to cover this configuration with the buildbots. Suggested-by: Tobias Grosser <tobias@grosser.es> llvm-svn: 298081
* [PruneUnprofitable] Add -polly-prune-unprofitable pass.Michael Kruse2017-03-171-0/+9
| | | | | | | | | | | | | | | | | | | | | | ScopInfo's normal profitability heuristic considers SCoPs where all statements have scalar writes as not profitably optimizable and invalidate the SCoP in that case. However, -polly-delicm and -polly-simplify may be able to remove some of the scalar writes such that the flag -polly-unprofitable-scalar-accs=false allows disabling that part of the heuristic. In cases where DeLICM (or other passes after ScopInfo) are not successful in removing scalar writes, the SCoP is still not profitably optimizable. The schedule optimizer would again try computing another schedule, resulting in slower compilation. The -polly-prune-unprofitable pass applies the profitability heuristic again before the schedule optimizer Polly can still bail out even with -polly-unprofitable-scalar-accs=false. Differential Revision: https://reviews.llvm.org/D31033 llvm-svn: 298080
* [ScopInfo] Introduce ScopStmt::getSurroundingLoop(). NFC.Michael Kruse2017-03-151-14/+0
| | | | | | | | | | | | | | Introduce ScopStmt::getSurroundingLoop() to replace getFirstNonBoxedLoopFor. getSurroundingLoop() returns the precomputed surrounding/first non-boxed loop. Except in ScopDetection, the list of boxed loops is only used to get the surrounding loop. getFirstNonBoxedLoopFor also requires LoopInfo at every use which is not necessarily available everywhere where we may want to use it. Differential Revision: https://reviews.llvm.org/D30985 llvm-svn: 297899
* [Simplify] Add -polly-simplify pass.Michael Kruse2017-03-101-0/+9
| | | | | | | | | | | | | | | | | This new pass removes unnecessary accesses and writes. It currently supports 2 simplifications, but more are planned. It removes write accesses that write a loaded value back to the location it was loaded from. It is a typical artifact from DeLICM. Removing it will get rid of bogus dependencies later in dependency analysis. It also removes statements without side-effects. ScopInfo already removes these, but the removal of unnecessary writes can result in more side-effect free statements. Differential Revision: https://reviews.llvm.org/D30820 llvm-svn: 297473
* Introduce isl C++ bindings, Part 1: value_ptr style interfaceTobias Grosser2017-03-102-114/+72
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Over the last couple of months several authors of independent isl C++ bindings worked together to jointly design an official set of isl C++ bindings which combines their experience in developing isl C++ bindings. The new bindings have been designed around a value pointer style interface and remove the need for explicit pointer managenent and instead use C++ language features to manage isl objects. This commit introduces the smart-pointer part of the isl C++ bindings and replaces the current IslPtr<T> classes, which served the very same purpose, but had to be manually maintained. Instead, we now rely on automatically generated classes for each isl object, which provide value_ptr semantics. An isl object has the following smart pointer interface: inline set manage(__isl_take isl_set *ptr); class set { friend inline set manage(__isl_take isl_set *ptr); isl_set *ptr = nullptr; inline explicit set(__isl_take isl_set *ptr); public: inline set(); inline set(const set &obj); inline set &operator=(set obj); inline ~set(); inline __isl_give isl_set *copy() const &; inline __isl_give isl_set *copy() && = delete; inline __isl_keep isl_set *get() const; inline __isl_give isl_set *release(); inline bool is_null() const; } The interface and behavior of the new value pointer style classes is inspired by http://www.open-std.org/jtc1/sc22/wg21/docs/papers/2012/n3339.pdf, which proposes a std::value_ptr, a smart pointer that applies value semantics to its pointee. We currently only provide a limited set of public constructors and instead require provide a global overloaded type constructor method "isl::obj isl::manage(isl_obj *)", which allows to convert an isl_set* to an isl::set by calling 'S = isl::manage(s)'. This pattern models the make_unique() constructor for unique pointers. The next two functions isl::obj::get() and isl::obj::release() are taken directly from the std::value_ptr proposal: S.get() extracts the raw pointer of the object managed by S. S.release() extracts the raw pointer of the object managed by S and sets the object in S to null. We additionally add std::obj::copy(). S.copy() returns a raw pointer refering to a copy of S, which is a shortcut for "isl::obj(oldobj).release()", a functionality commonly needed when interacting directly with the isl C interface where all methods marked with __isl_take require consumable raw pointers. S.is_null() checks if S manages a pointer or if the managed object is currently null. We add this function to provide a more explicit way to check if the pointer is empty compared to a direct conversion to bool. This commit also introduces a couple of polly-specific extensions that cover features currently not handled by the official isl C++ bindings draft, but which have been provided by IslPtr<T> and are consequently added to avoid code churn. These extensions include: - operator bool() : Conversion from objects to bool - construction from nullptr_t - get_ctx() method - take/keep/give methods, which match the currently used naming convention of IslPtr<T> in Polly. They just forward to (release/get/manage). - raw_ostream printers We expect that these extensions are over time either removed or upstreamed to the official isl bindings. We also export a couple of classes that have not yet been exported in isl (e.g., isl::space) As part of the code review, the following two questions were asked: - Why do we not use a standard smart pointer? std::value_ptr was a proposal that has not been accepted. It is consequently not available in the standard library. Even if it would be available, we want to expand this interface with a complete method interface that is conveniently available from each managed pointer. The most direct way to achieve this is to generate a specialiced value style pointer class for each isl object type and add any additional methods to this class. The relevant changes follow in subsequent commits. - Why do we not use templates or macros to avoid code duplication? It is certainly possible to use templates or macros, but as this code is auto-generated there is no need to make writing this code more efficient. Also, most of these classes will be specialized with individual member functions in subsequent commits, such that there will be little code reuse to exploit. Hence, we decided to do so at the moment. These bindings are not yet officially part of isl, but the draft is already very stable. The smart pointer interface itself did not change since serveral months. Adding this code to Polly is against our normal policy of only importing official isl code. In this case however, we make an exception to showcase a non-trivial use case of these bindings which should increase confidence in these bindings and will help upstreaming them to isl. Tags: #polly Reviewed By: Meinersbur Differential Revision: https://reviews.llvm.org/D30325 llvm-svn: 297452
* [Support] Add -polly-dump-module pass.Michael Kruse2017-03-092-0/+129
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | This pass allows writing the LLVM-IR just before and after the Polly passes to a file. Dumping the IR before Polly helps reproducing bugs that occur in code generated by clang. It is the only reliable way to get the IR that triggers a bug. The alternative is to emit the IR with clang -c -emit-llvm -S -o dump.ll then pass it through all optimization passes opt dump.ll -basicaa -sroa ... -S -o optdump.ll to then reproduce the error with opt optdump.ll -polly-opt-isl -polly-codegen -analyze However, the IR is not the same. -O3 uses a PassBuilder than creates passes with different parameters than the default. Dumping the IR after Polly is useful to compare a miscompilation with a known-good configuration. Differential Revision: https://reviews.llvm.org/D30788 llvm-svn: 297415
* Fix namespaces after clang-format updateTobias Grosser2017-03-011-1/+1
| | | | llvm-svn: 296635
* [Support] Remove NonowningIslPtr. NFC.Michael Kruse2017-02-232-22/+15
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | NonowningIslPtr<isl_X> was used as types of function parameters when the function does not consume the isl object, i.e. an __isl_keep parameter. The alternatives are: 1. IslPtr<isl_X> This has additional calls to isl_X_copy and isl_X_free to increase/decrease the reference counter even though not needed. The caller already owns a reference to the isl object. 2. const IslPtr<isl_X>& This does not change the reference counter, but requires an additional load to get the pointer to the isl object (instead of just passing the pointer itself). Moreover, the compiler cannot rely on the constness of the pointer and has to reload the pointer every time it writes to memory (unless alias analysis such as TBAA says it is not possible). The isl C++ bindings currently in development do not have an equivalent to NonowningIslPtr and adding one would make the binding more complicated and its advantage in performance is small. In order to simplify the transition to these C++ bindings, remove NonowningIslPtr. Change every former use of it to alternative 2 mentioned aboce (const IslPtr<isl_X>&). llvm-svn: 295998
* [Support] Add convertZoneToTimepoints. NFC.Michael Kruse2017-02-041-0/+16
| | | | | | | | | | | | This function has been extracted from the upcoming DeLICM patch (https://reviews.llvm.org/D24716). In contrast to computeReachingWrite and computeArrayUnused, convertZoneToTimepoints implies a format for zones (ranges between timepoints). Zones at the moment are unique to DeLICM, but convertZoneToTimepoints makes most sense in conjunction with the previous two functions. llvm-svn: 294094
* [Support] Add computeArrayUnused. NFC.Michael Kruse2017-02-041-0/+47
| | | | | | | This function has been extracted from the upcoming DeLICM patch (https://reviews.llvm.org/D24716). llvm-svn: 294093
* [Support] Add computeReachingWrite. NFC.Michael Kruse2017-02-041-0/+60
| | | | | | | This function has been extracted from the upcoming DeLICM patch (https://reviews.llvm.org/D24716). llvm-svn: 294092
* [Support] Remove unused function hasInvokeEdge. NFC.Michael Kruse2017-02-031-9/+0
| | | | llvm-svn: 294062
* Update to recent formatting changesTobias Grosser2017-02-013-14/+10
| | | | llvm-svn: 293756
* [Support] Add general isl tools for DeLICM. NFC.Michael Kruse2017-01-271-0/+264
| | | | | | | | | Add some generally useful isl tools into a their own new ISLTools.cpp. These are the helpers were extracted from and will be use by the DeLICM algorithm (https://reviews.llvm.org/D24716). Suggested-by: Tobias Grosser <tobias@grosser.es> llvm-svn: 293340
* Tidy up getFirstNonBoxedLoopFor [NFC]Eli Friedman2017-01-161-0/+14
| | | | | | | | | | | | Move the function getFirstNonBoxedLoopFor which is used in ScopBuilder and in ScopInfo to Support/ScopHelpers to make it reusable in other locations. No functionality change. Patch by Sameer Abu Asal. Differential Revision: https://reviews.llvm.org/D28754 llvm-svn: 292168
* Adjust formatting to commit r292110 [NFC]Tobias Grosser2017-01-162-4/+6
| | | | llvm-svn: 292123
* Add isl_multi_pw_aff to GICHelperAlexandre Isoard2016-12-161-0/+7
| | | | | | Add isl_multi_pw_aff* to GICHelper and add some missing isl_pw_multi_aff* handlers. llvm-svn: 290007
* Add more ISL foreachElt functions. NFC.Michael Kruse2016-12-071-0/+41
| | | | | | | Add and implement foreachElt for isl_map, isl_set and isl_union_set. These are used by an out-of-tree patch which is in process of being upstreamed. llvm-svn: 288924
* Add IslPtr type traits. NFC.Michael Kruse2016-12-071-1/+3
| | | | | | | | | | Add traits for isl_id and isl_multi_aff, required by out-of-tree patches currently in progress of upstreaming. isl_union_pw_aff_dump has been added to ISL during one of the last ISL updates, such that we can also enable its dump() trait. llvm-svn: 288915
* Allow to disable unsigned operations (zext, icmp ugt, ...)Johannes Doerfert2016-12-021-2/+20
| | | | | | | Unsigned operations are often useful to support but the heuristics are not yet tuned. This options allows to disable them if necessary. llvm-svn: 288521
* [FIX] Do not try to hoist obviously overwritten loadsJohannes Doerfert2016-12-011-0/+3
| | | | llvm-svn: 288328
* [DeLICM] Add pass boilerplate code.Michael Kruse2016-11-291-0/+10
| | | | | | | | | | Add an empty DeLICM pass, without any functional parts. Extracting the boilerplate from the the functional part reduces the size of the code to review (https://reviews.llvm.org/D24716) Suggested-by: Tobias Grosser <tobias@grosser.es> llvm-svn: 288160
* canSynthesize: Remove unused argument LI. NFC.Michael Kruse2016-11-291-2/+1
| | | | | | | The helper function polly::canSynthesize() does not directly use the LoopInfo analysis, hence remove it from its argument list. llvm-svn: 288144
* Probably overwritten loads should not be considered hoistableJohannes Doerfert2016-11-171-2/+22
| | | | | | | | Do not assume a load to be hoistable/invariant if the pointer is used by another instruction in the SCoP that might write to memory and that is always executed. llvm-svn: 287272
* [NFC] Add flag to disable error block assumptionsJohannes Doerfert2016-11-171-0/+7
| | | | | | | | | The declaration as an "error block" is currently aggressive and not very smart. This patch allows to disable error blocks completely. This might be useful to prevent SCoP expansion to a point where the assumed context becomes infeasible, thus the SCoP has to be discarded. llvm-svn: 287271
* [SCEVAffinator] Do not scan redundantly for parametersTobias Grosser2016-11-131-3/+0
| | | | | | | | | | | | | | | | | | | | In r286430 "SCEVValidator: add new parameters resulting from constant extraction" we added functionality to scan for parameters after constant extraction has taken place to ensure newly created parameters are correctly registered. This addition made the already existing registration of parameters redundant. Hence, we remove the corresponding call in this commit. An alternative solution would have been to also perform constant extraction when validating SCEV expressions and to then scan for parameters when validating a SCEV expression. However, as SCEV validation is used during SCoP detection where we want to be especially fast, adding additional functionality on this hot path should be avoided if good alternatives exist. In this case, we can choose to continue to only transform SCEV expression when actually modeling them. As all transformations we perform are expected to not change the validity of the SCEV expressions, this solution seems preferable. Suggested-by: Eli Friedman <efriedma@codeaurora.org> llvm-svn: 286780
* SCEVAffinator: pass parameter-only set to addRestriction if BB=nullptrTobias Grosser2016-11-101-0/+7
| | | | | | | | | | | | | | | | | | | | | | | | | | | Assumptions can either be added for a given basic block, in which case the set describing the assumptions is expected to match the dimensions of its domain. In case no basic block is provided a parameter-only set is expected to describe the assumption. The piecewise expressions that are generated by the SCEVAffinator sometimes have a zero-dimensional domain (e.g., [p] -> { [] : p <= -129 or p >= 128 }), which looks similar to a parameter-only domain, but is still a set domain. This change adds an assert that checks that we always pass parameter domains to addAssumptions if BB is empty to make mismatches here fail early. We also change visitTruncExpr to always convert to parameter sets, if BB is null. This change resolves http://llvm.org/PR30941 Another alternative to this change would have been to inspect all code to make sure we directly generate in the SCEV affinator parameter sets in case of empty domains. However, this would likely complicate the code which combines parameter and non-parameter domains when constructing a statement domain. We might still consider doing this at some point, but as this likely requires several non-local changes this should probably be done as a separate refactoring. Reported-by: Eli Friedman <efriedma@codeaurora.org> llvm-svn: 286444
* SCEVValidator: add new parameters resulting from constant extractionTobias Grosser2016-11-101-0/+3
| | | | | | | | | | | | | | | | | When extracting constant expressions out of SCEVs, new parameters may be introduced, which have not been registered before. This change scans SCEV expressions after constant extraction again to make sure newly introduced parameters are registered. We may for example extract the constant '8' from the expression '((8 * ((%a * %b) + %c)) + (-8 * %a))' and obtain the expression '(((-1 + %b) * %a) + %c)'. The new expression has a new parameter '(-1 + %b) * %a)', which was not registered before, but must be registered to not crash. This closes http://llvm.org/PR30953 Reported-by: Eli Friedman <efriedma@codeaurora.org> llvm-svn: 286430
* SCEVValidator: reduce indentation to increase readability [NFC]Tobias Grosser2016-11-081-12/+17
| | | | llvm-svn: 286217
* [Polly CodeGen] Break critical edge from RTC to original loop.Eli Friedman2016-11-021-8/+10
| | | | | | | | | | | | | | | This makes polly generate a CFG which is closer to what we want in LLVM IR, with a loop preheader for the original loop. This is just a cleanup, but it exposes some fragile assumptions. I'm not completely happy with the changes related to expandCodeFor; RTCBB->getTerminator() is basically a random insertion point which happens to work due to the way we generate runtime checks. I'm not sure what the right answer looks like, though. Differential Revision: https://reviews.llvm.org/D26053 llvm-svn: 285864
* [SCEVAffinator] Make precise modular math more correct.Eli Friedman2016-10-211-55/+25
| | | | | | | | | | | | | | | | | | | | | | | | | | | Integer math in LLVM IR is modular. Integer math in isl is arbitrary-precision. Modeling LLVM IR math correctly in isl requires either adding assumptions that math doesn't actually overflow, or explicitly wrapping the math. However, expressions with the "nsw" flag are special; we can pretend they're arbitrary-precision because it's undefined behavior if the result wraps. SCEV expressions based on IR instructions with an nsw flag also carry an nsw flag (roughly; actually, the real rule is a bit more complicated, but the details don't matter here). Before this patch, SCEV flags were also overloaded with an additional function: the ZExt code was mutating SCEV expressions as a hack to indicate to checkForWrapping that we don't need to add assumptions to the operand of a ZExt; it'll add explicit wrapping itself. This kind of works... the problem is that if anything else ever touches that SCEV expression, it'll get confused by the incorrect flags. Instead, with this patch, we make the decision about whether to explicitly wrap the math a bit earlier, basing the decision purely on the SCEV expression itself, and not its users. Differential Revision: https://reviews.llvm.org/D25287 llvm-svn: 284848
* Fix formatting after recent cl:: changesTobias Grosser2016-10-091-10/+13
| | | | | | This fixes 'make check-polly' llvm-svn: 283693
* Turn cl::values() (for enum) from a vararg function to using C++ variadic ↵Mehdi Amini2016-10-081-11/+7
| | | | | | | | | | | | | | | template The core of the change is supposed to be NFC, however it also fixes what I believe was an undefined behavior when calling: va_start(ValueArgs, Desc); with Desc being a StringRef. Differential Revision: https://reviews.llvm.org/D25342 llvm-svn: 283671
* [Support] Compile fix for gcc. NFC.Michael Kruse2016-09-301-2/+4
| | | | | | | | | | | | gcc 5.4 insists on template specialization to be in a namespace polly { ... } block, instead of being prefixed with 'polly::'. Error message: root/src/llvm/tools/polly/lib/Support/GICHelper.cpp:203:54: error: specialization of ‘template<class T> void polly::IslPtr<T>::dump() const’ in different namespace [-fpermissive] template <> void polly::IslPtr<isl_##TYPE>::dump() const { \ ^ msvc14 and clang 3.8 did not complain. llvm-svn: 282874
OpenPOWER on IntegriCloud