Commit Graph

3189 Commits

Author SHA1 Message Date
Siddharth Bhat 726c28f8c4 [CodeGen] Track trip counts per-scop for performance measurement.
- 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
2017-06-02 11:36:52 +00:00
Siddharth Bhat a4dea6bb05 [CodeGen] Print performance counter information in CSV.
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
2017-06-02 09:20:02 +00:00
Siddharth Bhat fee75f4ba5 [NFC] [CodeGen] Bail out of per-scop performance reporting if not supported.
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
2017-06-02 08:44:19 +00:00
Siddharth Bhat 07bee290de [CodeGen] Extend Performance Counter to track per-scop information.
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
2017-06-02 08:01:22 +00:00
Michael Kruse 3bb4829936 [CodeGen] Iterate over explicit instruction list for block statements. NFC
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
2017-06-02 00:13:49 +00:00
Michael Kruse 678aa336fa [ScopBuilder] Exclude ignored intrinsics from explicit instruction list.
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
2017-06-01 21:46:27 +00:00
Eli Friedman de1b318dad Add opt-bisect support to polly.
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
2017-06-01 21:29:05 +00:00
Tobias Grosser dff902fca7 [ScopInfo] Do not lookup key twice [NFC]
Suggested-by: Michael Kruse <llvm@meinersbur.de>
llvm-svn: 304410
2017-06-01 12:46:51 +00:00
Tobias Grosser f51decb5fe [BlockGenerator] Take context into account when identifying partial writes
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
2017-06-01 09:34:20 +00:00
Tobias Grosser 6b6ac90098 [BlockGenerator] Translate buildContainsCondition to idiomatic isl C++
llvm-svn: 304354
2017-05-31 21:49:51 +00:00
Tobias Grosser 5ecc5166d9 [isl++] Update bindings
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
2017-05-31 08:46:29 +00:00
Siddharth Bhat 0610b00295 [NFC] [PerfMonitor] Fix typo.
"InserBefore" -> "InsertBefore"

llvm-svn: 304287
2017-05-31 08:12:04 +00:00
Tobias Grosser 5863087ad3 [test] Add a short explanation to test
llvm-svn: 304279
2017-05-31 05:03:11 +00:00
Michael Kruse ed0c2f7e90 [ScopInfo] Do not add terminator & synthesizable instructions to the output 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
2017-05-29 12:27:38 +00:00
Siddharth Bhat 8bb436eb26 Revert "[NFC] Fix formatting & typecast issue. Build succeeds."
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
2017-05-29 11:34:29 +00:00
Siddharth Bhat ede801ca2b [NFC] Fix formatting & typecast issue. Build succeeds.
- 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
2017-05-29 11:00:31 +00:00
Tobias Grosser d9fb2842e7 Adapt to recent clang-format changes
llvm-svn: 304136
2017-05-29 08:06:29 +00:00
Tobias Grosser 1e55db30d5 Delinearize memory accesses that reference parameters coming from function calls
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
2017-05-27 15:18:53 +00:00
Tobias Grosser f5e7e60bc8 Allow side-effect free function calls in valid affine SCEVs
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
2017-05-27 15:18:46 +00:00
Tobias Grosser 7aa22859b6 Update some tests to changes in isl's internal representation
This was forgotten as part of r304069.

llvm-svn: 304070
2017-05-27 11:33:05 +00:00
Tobias Grosser 6ea64d8bd3 Update isl to isl-0.18-662-g17e172e
This is a general maintenance update

llvm-svn: 304069
2017-05-27 11:09:39 +00:00
Tobias Grosser 58822e8441 Update information on isl C++ bindings in Polly release notes
llvm-svn: 304068
2017-05-27 11:01:01 +00:00
Tobias Grosser d5fcbef8ee [Polly] Added the list of Instructions to output in ScopInfo pass
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
2017-05-27 04:40:18 +00:00
Tobias Grosser 9932086895 [ScopInfo] Translate mapToDimension to isl C++ [NFC]
llvm-svn: 304007
2017-05-26 17:22:03 +00:00
Siddharth Bhat c56a355a35 [Docs] Fix exception name being caught in case ReadTheDocs does not exist.
The exception is `ImportError`, not `ModuleNotFoundError`.

llvm-svn: 303982
2017-05-26 12:56:35 +00:00
Tobias Grosser 6e770813c9 Drop newline in docs builder to see if Polly docs are updated
llvm-svn: 303839
2017-05-25 05:38:05 +00:00
Tobias Grosser c8d13f50cc [ScopInfo] Tighten compute out introduced in r303404
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
2017-05-24 21:24:04 +00:00
Philip Pfaffe 1a0128faaa [Polly] Add handling of Top Level Regions
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
2017-05-24 18:39:39 +00:00
Michael Kruse 5f16986271 [DeLICM] Partial writes for PHIs.
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
2017-05-24 15:23:06 +00:00
Michael Kruse cb58bd6ccd [JSONImporter] misses checks whether the data it imports makes sense.
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
2017-05-24 15:09:35 +00:00
Tobias Grosser 8913015a79 Add a permanent Polly Hangouts URL
llvm-svn: 303725
2017-05-24 06:46:11 +00:00
Tobias Grosser 0bb785fc62 Add new C++ bindings to release notes
llvm-svn: 303724
2017-05-24 06:46:09 +00:00
Tobias Grosser 7eacc04836 Update version number in Polly release notes
llvm-svn: 303723
2017-05-24 06:46:06 +00:00
Tobias Grosser ecc0141349 Header file to help forcibly link GPURuntime
Summary: LinkGPURuntime.h defines and creates a structure ForceGPURuntimeLinking which creates an artificial dependency to functions defined in GPUJIT.c. The presence of this structure ensures that these functions are a part of the compiled object/library files including it.

Reviewers: grosser, Meinersbur

Reviewed By: grosser

Subscribers: #polly, pollydev

Tags: #polly

Differential Revision: https://reviews.llvm.org/D33198

llvm-svn: 303722
2017-05-24 06:46:03 +00:00
Siddharth Bhat 7bccbe780e [NFC] [PolyhedralInfo] fix typo
"an" -> "and"

llvm-svn: 303631
2017-05-23 11:41:52 +00:00
Philip Pfaffe 483340bb83 [Polly][NewPM] Reenable ScopPassManager unittest
llvm-svn: 303629
2017-05-23 11:28:50 +00:00
Philip Pfaffe 24a1bb2cf9 Post-commit fix of a comment
llvm-svn: 303628
2017-05-23 11:25:05 +00:00
Philip Pfaffe 78ae52f0d6 [Polly][NewPM] Port CodeGen to the new PM
Summary: To move CG to the new PM I outlined the various helper that were previously members of the CG class into free static functions. The CG class itself I moved into a header, which is required because we need to include it in `RegisterPasses` eventually.

Reviewers: grosser, Meinersbur

Reviewed By: grosser

Subscribers: pollydev, llvm-commits, sanjoy

Tags: #polly

Differential Revision: https://reviews.llvm.org/D33423

llvm-svn: 303624
2017-05-23 10:18:12 +00:00
Philip Pfaffe 2b852e2e42 [Polly][NewPM] Port IslAst to the new ScopPassManager
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
2017-05-23 10:12:56 +00:00
Philip Pfaffe 3ef36fa222 [Polly][NewPM] Port DependenceInfo to the new ScopPassManager.
Summary: This patch ports DependenceInfo to the new ScopPassManager. Printing is implemented as a seperate printer pass.

Reviewers: grosser, Meinersbur

Reviewed By: grosser

Subscribers: llvm-commits, pollydev

Tags: #polly

Differential Revision: https://reviews.llvm.org/D33421

llvm-svn: 303621
2017-05-23 10:09:06 +00:00
Tobias Grosser a32de1341e [ScopInfo] Translate foldAccessRelation to isl C++ [NFC]
llvm-svn: 303615
2017-05-23 07:22:56 +00:00
Tobias Grosser 53fc355e7d [ScopInfo] Translate buildMemIntrinsicAccessRelation to isl C++ [NFC]
llvm-svn: 303612
2017-05-23 07:07:09 +00:00
Tobias Grosser 1e2edaf3ea [ScopInfo] Translate assumeNoOutOfBound to isl C++ [NFC]
llvm-svn: 303611
2017-05-23 07:07:07 +00:00
Tobias Grosser b1ed3d9749 [ScopInfo] Translate applyAndSetFAD to isl C++
llvm-svn: 303610
2017-05-23 07:07:05 +00:00
Tobias Grosser 2ade986c9e [ScopInfo] Translate isReadOnly to isl C++
llvm-svn: 303608
2017-05-23 06:41:04 +00:00
Tobias Grosser 6d459c5d3d [ScopInfo] Simplify domains early
This speeds up scop modeling for scops with many redundent existentially
quantified constraints. For the attached test case, this change reduces
scop modeling time from minutes (hours?) to 0.15 seconds.

This change resolves a compilation timeout on the AOSP build.

Thanks Eli for reporting _and_ reducing the test case!

Reported-by: Eli Friedman <efriedma@codeaurora.org>
llvm-svn: 303600
2017-05-23 04:26:28 +00:00
Michael Kruse 1aad76c18f [CodeGen] Add invalidation of the loop SCEVs after merge block generation.
The SCEVs of loops surrounding the escape users of a merge blocks are
forgotten, so that loop trip counts based on old values can be revoked.

This fixes llvm.org//PR32536

Contributed-by: Baranidharan Mohan <mbdharan@gmail.com>

Differential Revision: https://reviews.llvm.org/D33195

llvm-svn: 303561
2017-05-22 15:36:53 +00:00
Siddharth Bhat 68ae83e68c [Docs] Use ReadTheDocs theme if available.
Use ReadTheDocs theme for Sphinx if available since it is well
maintained and used by readthedocs.org.

Differential Revision: https://reviews.llvm.org/D33387

llvm-svn: 303550
2017-05-22 13:36:15 +00:00
Siddharth Bhat b2f754e39f [Docs] Fix Sphinx documentation in CMake check.
Summary:
- `include(AddSphinxTarget)` needs to occur before checking `SPHINX_FOUND`.
- `docs-polly-html` and `docs-polly-man` are now usable again.
- Perhaps we should build docs in the CI as well?

Differential Revision: https://reviews.llvm.org/D33386

llvm-svn: 303549
2017-05-22 13:16:02 +00:00
Michael Kruse 706f79ab14 [CodeGen] Support partial write accesses.
Allow the BlockGenerator to generate memory writes that are not defined
over the complete statement domain, but only over a subset of it. It
generates a condition that evaluates to 1 if executing the subdomain,
and only then execute the access.

Only write accesses are supported. Read accesses would require a PHINode
which has a value if the access is not executed.

Partial write makes DeLICM able to apply mappings that are not defined
over the entire domain (for instance, a branch that leaves a loop with
a PHINode in its header; a MemoryKind::PHI write when leaving is never
read by its PHI read).

Differential Revision: https://reviews.llvm.org/D33255

llvm-svn: 303517
2017-05-21 22:46:57 +00:00