summaryrefslogtreecommitdiffstats
path: root/polly
Commit message (Collapse)AuthorAgeFilesLines
...
* [Docs] Add a catch-all file to document workflow & tricks.Siddharth Bhat2017-06-152-0/+40
| | | | | | | | - Write about `bugpoint` and `git-svn` for now. Differential Revision: https://reviews.llvm.org/D34147 llvm-svn: 305464
* Don't check side effects for functions outside of SCoPEli Friedman2017-06-142-4/+41
| | | | | | | | | | | | In r304074 we introduce a patch to accept results from side effect free functions into SCEV modeling. This causes rejection of cases where the call is happening outside the SCoP. This patch checks if the call is outside the Region and treats the results as a parameter (SCEVType::PARAM) to the SCoP instead of returning SCEVType::INVALID. Patch by Sameer Abu Asal. llvm-svn: 305423
* [www] Add missing '.' at the end of the sentenceTobias Grosser2017-06-131-1/+1
| | | | llvm-svn: 305336
* [www] Add the new AOSP buildbot to release notesTobias Grosser2017-06-131-0/+14
| | | | llvm-svn: 305334
* [www] Remove outdated documentationTobias Grosser2017-06-123-617/+0
| | | | | | | | | | Remove examples 'load_Polly_into_clang' and 'manual_matmul'. This information is now available in our SPHINX docs (*). (*) Thanks to Singapuram Sanjay Srivallabh <singapuram.sanjay@gmail.com> who contributed the SPHINX docs update! llvm-svn: 305186
* [Polly] [PPCGCodeGeneration] Skip Scops which contain function pointers.Siddharth Bhat2017-06-122-0/+120
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | In `PPCGCodeGeneration`, we try to take the references of every `Value` that is used within a Scop to offload to the kernel. This occurs in `GPUNodeBuilder::createLaunchParameters`. This breaks if one of the values is a function pointer, since one of these cases will trigger: 1. We try to to take the references of an intrinsic function, and this breaks at `verifyModule`, since it is illegal to take the reference of an intrinsic. 2. We manage to take the reference to a function, but this fails at `verifyModule` since the function will not be present in the module that is created in the kernel. 3. Even if `verifyModule` succeeds (which should not occur), we would then try to call a *host function* from the *device*, which is illegal runtime behaviour. So, we disable this entire range of possibilities by simply not allowing function references within a `Scop` which corresponds to a kernel. However, note that this is too conservative. We *can* allow intrinsics within kernels if the backend can lower the intrinsic correctly. For example, an intrinsic like `llvm.powi.*` can actually be lowered by the `NVPTX` backend. We will now gradually whitelist intrinsics which are known to be safe. Differential Revision: https://reviews.llvm.org/D33414 llvm-svn: 305185
* [NFC] Fix typo in `ImportJScop` declaration.Siddharth Bhat2017-06-121-1/+1
| | | | | | | | Contributed by: Singapuram Sanjay Differential Revision: https://reviews.llvm.org/D34079 llvm-svn: 305183
* [isl-cpp] Remove isl/mat.h and add insert_partial_scheduleTobias Grosser2017-06-111-40/+22
| | | | | | | | | | The isl/mat.h functionality was incomplete (we returned 'void *' instead of 'isl::mat') and is likely not needed. *.insert_partial_schedule was until know not exported in the bindings, but will be needed in the next step. llvm-svn: 305161
* [Polly] [ScopDetection] Allow passing multiple functions to `-polly-only-func`.Siddharth Bhat2017-06-092-5/+113
| | | | | | | | - This is useful to run optimisations on only certain functions. Differential Revision: https://reviews.llvm.org/D33990 llvm-svn: 305060
* Fix a lot of typos. NFC.Michael Kruse2017-06-0829-83/+83
| | | | llvm-svn: 304974
* [ScopInfo] Translate getNonHoistableCtx to C++ [NFC]Tobias Grosser2017-06-062-38/+25
| | | | llvm-svn: 304841
* [CMake] Retire Polly's FindCUDA.cmake in favour of CMake's default ↵Philip Pfaffe2017-06-064-29/+6
| | | | | | | | | | | | | | | | | | FindCUDA.cmake script. Patch by: Singapuram Sanjay Reviewers: grosser, philip.pfaffe Reviewed By: philip.pfaffe Subscribers: pollydev, mgorny Tags: #polly Differential Revision: https://reviews.llvm.org/D33669 llvm-svn: 304818
* [JScop] Emit error messages on error.Michael Kruse2017-06-061-3/+9
| | | | | | In importArrays instead of silently ignoring the file. llvm-svn: 304817
* [Simplify] Use execution order of memory accesses.Michael Kruse2017-06-064-1/+128
| | | | | | | | | | | | Iterate through memory accesses in execution order (first all implicit reads, then explicit accesses, then implicit writes). In the test case this caused an implicit load to be handled as if it was loaded after the write. That is, the value being written before it is available. This fixes llvm.org/PR33323 llvm-svn: 304810
* [Polly] [BlockGen] Support partial writes in regionsTobias Grosser2017-06-068-38/+400
| | | | | | | | | | | | | | | | | | | | | | | | Summary: The RegionGenerator traditionally kept a BlockMap that mapped from original basic blocks to newly generated basic blocks. With the introduction of partial writes such a 1:1 mapping is not possible any more, as a single basic block can be code generated into multiple basic blocks. Hence, depending on the use case we need to either use the first basic block or the last basic block. This is intended to address the last four cases of incorrect code generation in our AOSP buildbot and hopefully should turn it green. Reviewers: Meinersbur, bollu, gareevroman, efriedma, huihuiz, sebpop, simbuerg Reviewed By: Meinersbur Subscribers: pollydev, llvm-commits Tags: #polly Differential Revision: https://reviews.llvm.org/D33767 llvm-svn: 304808
* [CodeGen] Remove extra ';'. NFC.Michael Kruse2017-06-061-1/+1
| | | | | | | | | | Fix compiler warning: polly/lib/CodeGen/PerfMonitor.cpp:81:2: warning: extra ‘;’ [-Wpedantic] }; ^ llvm-svn: 304802
* [www] Fix links to bug trackerTobias Grosser2017-06-061-11/+10
| | | | | | | | | | | The LLVM bug tracker is now available at bugs.llvm.org instead of llvm.org/bugs. By updating our links to the tracker we do not only avoid unnecessary redirects, but also certificate warnings. We use this opportunity to shorten the text and to rename the link 'open bugs' to 'show open bugs' to clearify its meaning. llvm-svn: 304768
* Update isl to isl-0.18-679-g6e75a0dTobias Grosser2017-06-0412-152/+771
| | | | | | This is a regular maintenance update llvm-svn: 304686
* Add test coverage for regions with non-affine loopsTobias Grosser2017-06-031-0/+62
| | | | | | | | This adds test coverage for regions with non-affine loops, which we unfortunately missed when committing this features years ago. We will add more test coverage over time. llvm-svn: 304672
* [CodeGen] Track trip counts per-scop for performance measurement.Siddharth Bhat2017-06-025-104/+173
| | | | | | | | | | | | | | | | | | | - Add a counter that is incremented once on exit from a scop. - Test cases got split into two: one to test the cycles, and another one to test trip counts. - Sample output: ```name=sample-output.txt scop function, entry block name, exit block name, total time, trip count warmup, %entry.split, %polly.merge_new_and_old, 5180, 1 f, %entry.split, %polly.merge_new_and_old, 409944, 500 g, %entry.split, %polly.merge_new_and_old, 1226, 1 ``` Differential Revision: https://reviews.llvm.org/D33822 llvm-svn: 304543
* [CodeGen] Print performance counter information in CSV.Siddharth Bhat2017-06-022-10/+18
| | | | | | | | | | | | | | | | | | | | | | | | This ensures that tools can parse performance information which Polly generates easily. - Sample output: ```name=out.csv scop function, entry block name, exit block name, total time warmup, %entry.split, %polly.merge_new_and_old, 1960 f, %entry.split, %polly.merge_new_and_old, 1238 g, %entry.split, %polly.merge_new_and_old, 1218 ``` - Example code to parse output: ```lang=python, name=example-parse.py import asciitable import sys table = asciitable.read('out.csv', delimiter=',') asciitable.write(table, sys.stdout, delimiter=',') ``` llvm-svn: 304533
* [NFC] [CodeGen] Bail out of per-scop performance reporting if not supported.Siddharth Bhat2017-06-021-0/+8
| | | | | | | | | | We should bail out if performance monitoring is not supported, since we would have no information to print per-scop, and `FinalStartBB`, `ReturnFromFinal` would be `nullptr`. Assert that these are not `nullptr` if performance monitoring is supported. llvm-svn: 304529
* [CodeGen] Extend Performance Counter to track per-scop information.Siddharth Bhat2017-06-027-33/+211
| | | | | | | | | | | | | | | Previously, we would generate one performance counter for all scops. Now, we generate both the old information, as well as a per-scop performance counter to generate finer grained information. This patch needed a way to generate a unique name for a `Scop`. The start region, end region, and function name combined provides a unique `Scop` name. So, `Scop` has a new public API to provide its start and end region names. Differential Revision: https://reviews.llvm.org/D33723 llvm-svn: 304528
* [CodeGen] Iterate over explicit instruction list for block statements. NFCMichael Kruse2017-06-022-2/+10
| | | | | | | | | | | | For when statements do not contain all instructions of a BasicBlock anymore, the block generator needs to go through the explicit list of instructions it contains. Contributed-by: Nandini Singhal <cs15mtech01004@iith.ac.in> Differential Revision: https://reviews.llvm.org/D33653 llvm-svn: 304502
* [ScopBuilder] Exclude ignored intrinsics from explicit instruction list.Michael Kruse2017-06-012-1/+48
| | | | | | | | | | | | | | Ignored intrinsics are ignored at code generation, therefore do not need to be part of the instruction list. Specifically, llvm.lifetime.* intrinisics are removed before code generation, referencing them would cause a use-after-free error. Contributed-by: Nandini Singhal <cs15mtech01004@iith.ac.in> Differential Revision: https://reviews.llvm.org/D33768 llvm-svn: 304483
* Add opt-bisect support to polly.Eli Friedman2017-06-012-1/+7
| | | | | | | | | This is useful for debugging miscompiles and extracting testcases for crashes. See http://llvm.org/docs/OptBisect.html . Differential Revision: https://reviews.llvm.org/D33752 llvm-svn: 304480
* [ScopInfo] Do not lookup key twice [NFC]Tobias Grosser2017-06-011-1/+1
| | | | | Suggested-by: Michael Kruse <llvm@meinersbur.de> llvm-svn: 304410
* [BlockGenerator] Take context into account when identifying partial writesTobias Grosser2017-06-012-1/+38
| | | | | | | | | | | | | | | | | | | | | | | | | | A partial write is a write where the domain of the values written is a subset of the execution domain of the parent statement containing the write. Originally, we directly checked this subset relation whereas it is indeed only important that the subset relation holds for the parameter values that are known to be valid in the execution context of the scop. We update our check to avoid the unnecessary introduction of partial writes in situations where the write appears to be partial without context information, but where context information allows us to understand that a full write can be generated. This change fixes (hides) a recent regression introduced in r303517, which broke our AOSP builds. The part that is correctly fixed in this change is that we do not any more unnecessarily generate a partial write. This is good performance wise and, as we currently do not yet explicitly introduce partial writes in the default configuration, this also hides possible bugs in the partial writes implementation. The crashes that we have originally seen were caused by such a bug, where partial writes were incorrectly generated in region statements. An additional patch in a subsequent commit is needed to address this problem. Reported-by: Reported-by: Eli Friedman <efriedma@codeaurora.org> Differential Revision: https://reviews.llvm.org/D33759 llvm-svn: 304398
* [BlockGenerator] Translate buildContainsCondition to idiomatic isl C++Tobias Grosser2017-05-311-18/+11
| | | | llvm-svn: 304354
* [isl++] Update bindingsTobias Grosser2017-05-311-1/+13
| | | | | | | | This change removes the requirement for explicit conversions from isl::boolean to isl::bool, which resolves a compilation error on OSX. Suggested-by: Siddharth Bhat <siddu.druid@gmail.com> llvm-svn: 304288
* [NFC] [PerfMonitor] Fix typo.Siddharth Bhat2017-05-311-1/+1
| | | | | | "InserBefore" -> "InsertBefore" llvm-svn: 304287
* [test] Add a short explanation to testTobias Grosser2017-05-311-0/+6
| | | | llvm-svn: 304279
* [ScopInfo] Do not add terminator & synthesizable instructions to the output ↵Michael Kruse2017-05-292-7/+5
| | | | | | | | | | | instructions. Such instructions are generates on-demand by the CodeGenerator and thus do not need representation in a statement. Differential Revision: https://reviews.llvm.org/D33642 llvm-svn: 304151
* Revert "[NFC] Fix formatting & typecast issue. Build succeeds."Siddharth Bhat2017-05-292-2/+2
| | | | | | | | Should not have 'fixed' the formatting issue, I did not have the most recent version of `clang-format`. This reverts commit 761b1268359e14e59142f253d77864a29d55c56c. llvm-svn: 304148
* [NFC] Fix formatting & typecast issue. Build succeeds.Siddharth Bhat2017-05-292-2/+2
| | | | | | | | | - Fix formatting in `RegisterPasses.cpp`. - `assert` tried to compare `isl::boolean` against `long`. Explicitly construct `bool` from `isl::boolean`. This allows the implicit cast of `bool` to `long. llvm-svn: 304146
* Adapt to recent clang-format changesTobias Grosser2017-05-291-1/+1
| | | | llvm-svn: 304136
* Delinearize memory accesses that reference parameters coming from function callsTobias Grosser2017-05-275-5/+188
| | | | | | | | | | | | | | 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-274-3/+143
| | | | | | | | | | | | | | | 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
* Update some tests to changes in isl's internal representationTobias Grosser2017-05-275-9/+11
| | | | | | This was forgotten as part of r304069. llvm-svn: 304070
* Update isl to isl-0.18-662-g17e172eTobias Grosser2017-05-2767-812/+2673
| | | | | | This is a general maintenance update llvm-svn: 304069
* Update information on isl C++ bindings in Polly release notesTobias Grosser2017-05-271-9/+8
| | | | llvm-svn: 304068
* [Polly] Added the list of Instructions to output in ScopInfo passTobias Grosser2017-05-274-7/+89
| | | | | | | | | | | | Summary: This patch outputs all the list of instructions in BlockStmts. Reviewers: Meinersbur, grosser, bollu Subscribers: bollu, llvm-commits, pollydev Differential Revision: https://reviews.llvm.org/D33163 llvm-svn: 304062
* [ScopInfo] Translate mapToDimension to isl C++ [NFC]Tobias Grosser2017-05-261-46/+18
| | | | llvm-svn: 304007
* [Docs] Fix exception name being caught in case ReadTheDocs does not exist.Siddharth Bhat2017-05-261-1/+1
| | | | | | The exception is `ImportError`, not `ModuleNotFoundError`. llvm-svn: 303982
* Drop newline in docs builder to see if Polly docs are updatedTobias Grosser2017-05-251-1/+0
| | | | llvm-svn: 303839
* [ScopInfo] Tighten compute out introduced in r303404Tobias Grosser2017-05-241-1/+1
| | | | | | | It seems we are still spending too much time on rare inputs, which continue to timeout the AOSP buildbot. Let's see if a further reduction is sufficient. llvm-svn: 303807
* [Polly] Add handling of Top Level RegionsPhilip Pfaffe2017-05-245-16/+34
| | | | | | | | | | | | | | | | | | | | | | | | | | | 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
* [DeLICM] Partial writes for PHIs.Michael Kruse2017-05-244-1/+173
| | | | | | | | | | | | | | Enable the use for partial writes for PHI write accesses with a switch. This simply skips the test for whether a PHI write would be partial. The analog test for partial value writes also protects for partial reads which we do not support (yet). It is possible to test for partial reads separately such that we could skip the partial write check as well. In case this shows up to be useful, I can implement it as well. Differential Revision: https://reviews.llvm.org/D33487 llvm-svn: 303762
* [JSONImporter] misses checks whether the data it imports makes sense.Michael Kruse2017-05-2441-7/+2084
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Without this patch, the JSONImporter did not verify if the data it loads were correct or not (Bug llvm.org/PR32543). I add some checks in the JSONImporter class and some test cases. Here are the checks (and test cases) I added : JSONImporter::importContext - The "context" key does not exist. - The context was not parsed successfully by ISL. - The isl_set has the wrong number of parameters. - The isl_set is not a parameter set. JSONImporter::importSchedule - The "statements" key does not exist. - There is not the right number of statement in the file. - The "schedule" key does not exist. - The schedule was not parsed successfully by ISL. JSONImporter::importAccesses - The "statements" key does not exist. - There is not the right number of statement in the file. - The "accesses" key does not exist. - There is not the right number of memory accesses in the file. - The "relation" key does not exist. - The memory access was not parsed successfully by ISL. JSONImporter::areArraysEqual - The "type" key does not exist. - The "sizes" key does not exist. - The "name" key does not exist. JSONImporter::importArrays /!\ Do not check if there is an key name "arrays" because it is not considered as an error. All checks are already in place or implemented in JSONImporter::areArraysEqual. Contributed-by: Nicolas Bonfante <nicolas.bonfante@insa-lyon.fr> Differential Revision: https://reviews.llvm.org/D32739 llvm-svn: 303759
* Add a permanent Polly Hangouts URLTobias Grosser2017-05-241-0/+15
| | | | llvm-svn: 303725
OpenPOWER on IntegriCloud