Commit Graph

1127 Commits

Author SHA1 Message Date
Eli Friedman 127e0cd21b Don't check side effects for functions outside of SCoP
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
2017-06-14 22:43:28 +00:00
Siddharth Bhat bccaea57c0 [Polly] [PPCGCodeGeneration] Skip Scops which contain function pointers.
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
2017-06-12 11:41:09 +00:00
Siddharth Bhat 286c916dde [Polly] [ScopDetection] Allow passing multiple functions to `-polly-only-func`.
- This is useful to run optimisations on only certain functions.

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

llvm-svn: 305060
2017-06-09 08:23:40 +00:00
Michael Kruse ad7a1805be [Simplify] Use execution order of memory accesses.
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
2017-06-06 17:46:42 +00:00
Tobias Grosser deefbced96 [Polly] [BlockGen] Support partial writes in regions
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
2017-06-06 17:17:30 +00:00
Tobias Grosser 22be8a18f3 Add test coverage for regions with non-affine loops
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
2017-06-03 23:39:02 +00:00
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 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 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
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 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
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 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
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 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
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
Tobias Grosser ee61ebb134 Fix buildbots after r303429
A test case with a GPU runline was added without setting 'REQUIRES=pollyacc'. We
drop the GPU run line, as the basic functionality can already be tested with
the normal code generation.

llvm-svn: 303485
2017-05-20 04:22:26 +00:00
Siddharth Bhat b7f68b8c9e [Fortran Support] Materialize outermost dimension for Fortran array.
- We use the outermost dimension of arrays since we need this
information to generate GPU transfers.

- In general, if we do not know the outermost dimension of the array
(because the indexing expression is non-affine, for example) then we
simply cannot generate transfer code.

- However, for Fortran arrays, we can use the Fortran array
representation which stores the dimensions of all arrays.

- This patch uses the Fortran array representation to generate code that
computes the outermost dimension size.

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

llvm-svn: 303429
2017-05-19 15:07:45 +00:00
Tobias Grosser d8945baa0a [ScopDetection] Allow detection of full functions
This is useful when only analyzing functions.

llvm-svn: 303420
2017-05-19 12:13:02 +00:00
Tobias Grosser 45e9fd1810 [ScopInfo] Gracefully handle long compile times
The following test case tried to compute the lexicographic minimum of the
following set during alias analysis, which caused very long compile time:

[p_0, p_1, p_2, p_3, p_4, p_5] -> { MemRef0[i0] : (517p_3 >= 70944 - 298p_2 and
256i0 >= -71199 + 298p_2 + 517p_3 and 256i0 <= -70944 + 298p_2 + 517p_3) or
(409p_4 >= 57120 - 298p_2 and 256i0 >= -57375 + 298p_2 + 409p_4 and 256i0 <=
-57120 + 298p_2 + 409p_4) or (104p_4 >= 17329 + 149p_2 - 50p_3 and 128i0 >=
17328 + 149p_2 - 50p_3 - 104p_4 and 128i0 <= 17455 + 149p_2 - 50p_3 - 104p_4) or
(104p_4 <= 17328 + 149p_2 - 50p_3 and 128i0 >= 17201 + 149p_2 - 50p_3 - 104p_4
and 128i0 <= 17328 + 149p_2 - 50p_3 - 104p_4) or (409p_4 <= 57119 - 298p_2 and
256i0 >= -57120 + 298p_2 + 409p_4 and 256i0 <= -56865 + 298p_2 + 409p_4) or
(517p_3 <= 70943 - 298p_2 and 256i0 >= -70944 + 298p_2 + 517p_3 and 256i0 <=
-70689 + 298p_2 + 517p_3) or (p_1 >= 2 + 2p_0 and 298p_5 >= 70944 - 517p_3 and
256i0 >= -71199 + 517p_3 + 298p_5 and 256i0 <= -70944 + 517p_3 + 298p_5) or (p_1
>= 2 + 2p_0 and 298p_5 >= 57120 - 409p_4 and 256i0 >= -57375 + 409p_4 + 298p_5
>and 256i0 <= -57120 + 409p_4 + 298p_5) or (p_1 >= 2 + 2p_0 and 149p_5 <= -17329
>+ 50p_3 + 104p_4 and 128i0 >= 17328 - 50p_3 - 104p_4 + 149p_5 and 128i0 <=
>17455 - 50p_3 - 104p_4 + 149p_5) or (p_1 >= 2 + 2p_0 and 149p_5 >= -17328 +
>50p_3 + 104p_4 and 128i0 >= 17201 - 50p_3 - 104p_4 + 149p_5 and 128i0 <= 17328
>- 50p_3 - 104p_4 + 149p_5) or (p_1 >= 2 + 2p_0 and 298p_5 <= 57119 - 409p_4 and
>256i0 >= -57120 + 409p_4 + 298p_5 and 256i0 <= -56865 + 409p_4 + 298p_5) or
>(p_1 >= 2 + 2p_0 and 298p_5 <= 70943 - 517p_3 and 256i0 >= -70944 + 517p_3 +
>298p_5 and 256i0 <= -70689 + 517p_3 + 298p_5) }

We now guard the potentially expensive functions in Polly's scop analysis to
gracefully bail out in case of overly long compilation times.

llvm-svn: 303404
2017-05-19 03:45:00 +00:00
Siddharth Bhat 06e3c74d83 [Fortran Support] Change "global" pattern match to work for params
Summary:
- Rename global / local naming convention that did not make much sense
to Visible / Invisible, where the visible refers to whether the ALLOCATE
call to the Fortran array is present in the current module or not.

- This match now works on both cross fortran module globals and on
parameters to functions since neither of them are necessarily allocated
at the point of their usage.

- Add testcase that matches against both a load and a store against
function parameters.

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

llvm-svn: 303356
2017-05-18 16:47:13 +00:00
Philip Pfaffe 3030bf0c81 [Polly][Fortran Support] Fix two testcases for the loadable-library use-case
llvm-svn: 303057
2017-05-15 12:58:31 +00:00
Siddharth Bhat 0fe7231a2f [Fortran Support] Add pattern match for Fortran Arrays that are parameters.
- This breaks the previous assumption that Fortran Arrays are `GlobalValue`.

- The names of functions were getting unwieldy. So, I renamed the
Fortran related functions.

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

llvm-svn: 303040
2017-05-15 08:41:30 +00:00
Tobias Grosser b693f42b71 [Polly] Fix code generation of llvm.expect intrinsic
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
2017-05-14 09:09:54 +00:00
Michael Kruse fa7be88378 [Simplify] Remove identical write removal. NFC.
Removal of overwritten writes currently encompasses all the cases
of the identical write removal.

There is an observable behavioral change in that the last, instead
of the first, MemoryAccess is kept. This should not affect the
generated code, however.

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

llvm-svn: 302987
2017-05-13 12:20:57 +00:00
Michael Kruse f263610b82 [Simplify] Remove writes that are overwritten.
Remove memory writes that are overwritten by later writes. This works
for StoreInsts:

      store double 21.0, double* %A
      store double 42.0, double* %A

scalar writes at the end of a statement and mixes of these.

Multiple writes can be the result of DeLICM, which might map multiple
writes to the same location when it knows that these do no conflict
(for instance because they write the same value). Such writes
interfere with pattern-matched optimization such as gemm and may not
get removed by other LLVM passes after code generation.

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

llvm-svn: 302986
2017-05-13 11:49:34 +00:00
Siddharth Bhat d0d29addf9 [NFC] [Fortran Support] Run -instnamer on testcases
llvm-svn: 302892
2017-05-12 12:36:04 +00:00
Siddharth Bhat f16db04cd5 [FIX] Fix regression caused by c29f4ed, testcase matches output
- Commit changed codegen for induction variables
- Updated testcase

llvm-svn: 302891
2017-05-12 11:34:51 +00:00
Siddharth Bhat c05fcc0d9e [NFC] [Fortran Support] Cleanup Fortran Array pattern mactch testcases
- Move the testcases to ScopInfo/ since the processing takes place in
  ScopBuilder.

- Cleanup testcases, run -polly-canonicalize on them, find minimal set
  of opt parameters.

llvm-svn: 302886
2017-05-12 09:37:39 +00:00
Hongbin Zheng 4fe342cb75 [Polly] Generate more 'canonical' induction variable
Today Polly generates induction variable in this way:

polly.indvar = phi 0, polly.indvar.next
...
polly.indvar.next = polly.indvar + stide
polly.loop_cond = predicate polly.indvar, (UB - stride)

Instead of:

polly.indvar = phi 0, polly.indvar.next
...
polly.indvar.next = polly.indvar + stide
polly.loop_cond = predicate polly.indvar.next, UB

The way Polly generate induction variable cause some problem in the indvar simplify pass.
This patch make polly generate the later form, by assuming the induction variable never overflow

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

llvm-svn: 302866
2017-05-12 02:17:15 +00:00
Michael Kruse 07e315e780 [Simplify] Remove identical scalar writes.
After DeLICM, it is possible to have two writes of the same value to
the same location in the same statement when it determined that those
writes do not conflict (write the same value).

Teach -polly-simplify to remove one of the writes. It interferes with
the pattern matching of matrix-multiplication kernels and also seem
to not be optimized away by LLVM.

The algorthm is simple, has O(n^2) behaviour (n = max number of
MemoryAccesses in a statement) and only matches the most obvious cases,
but seem to be enough to pattern-match Boost ublas gemm.

Not handled cases include:
- StoreInst instructions (a.k.a. explicit writes), since the value might
  be loaded or overwritten between the two stores.
- PHINode, especially LCSSA, when the PHI value matches with on other's.
- Partial writes (in preparation)

llvm-svn: 302805
2017-05-11 15:07:38 +00:00
Siddharth Bhat abea18feba [NFC] [Fortran Support] move Fortran array detection testcases
move these testcases to where they belong: ScopDetect

llvm-svn: 302735
2017-05-10 21:35:14 +00:00
Siddharth Bhat f5c81fb199 [Fix][Fortran Support] Don't use -debug-only in pattern matching test cases
-debug-only is unnecessary and causes the tests to break in Release
mode. Remove the option to opt in the test cases.

llvm-svn: 302722
2017-05-10 20:10:17 +00:00
Michael Kruse f69a7c306b [DeLICM] Always normalize domain. NFC.
Some isl functions can simplify their __isl_keep arguments. The
argument object after the call uses different contraints to represent
the same set. Different contraints can result in different outputs
when printed to a string.

In assert builds additional isl functions are called (in assert() or
mentioned, these can change the internal representation of its read-only
arguments such that printed strings are different in debug and non-debug
builds.

What happened here is that a call to isl_set_is_equal inside an assert
in getScatterFor normalizes one of its arguments such that one redundant
constraint is removed. The redundant constraint therefore does not appear
in the string representing the domain, which FileCheck notices as a
regression test failure compared to a build with assertions disabled.

This fix removes the redundant contraints the domain from the start such
that the redundant contraint is removed in assert and non-assert builds.
Isl adds a flag to such sets such that the removal of redundancies is
not done multiple times (here: by isl_set_is_equal).

Thanks to Tobias Grosser for reporting and hinting to the cause.

llvm-svn: 302711
2017-05-10 19:50:45 +00:00
Siddharth Bhat f2dbba8183 [Fortran Support] Detect Fortran arrays & metadata from dragonegg output
Add the ability to tag certain memory accesses as those belonging to
Fortran arrays. We do this by pattern matching against known patterns
of Dragonegg's LLVM IR output from Fortran code.

Fortran arrays have metadata stored with them in a struct. This struct
is called the "Fortran array descriptor", and a reference to this is
stored in each MemoryAccess.

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

llvm-svn: 302653
2017-05-10 13:11:20 +00:00
Tobias Grosser f3adab4c20 [Polly] Canonicalize arrays according to base-ptr equivalence class
Summary:
    In case two arrays share base pointers in the same invariant load equivalence
    class, we canonicalize all memory accesses to the first of these arrays
    (according to their order in the equivalence class).

    This enables us to optimize kernels such as boost::ublas by ensuring that
    different references to the C array are interpreted as accesses to the same
    array. Before this change the runtime alias check for ublas would fail, as it
    would assume models of the C array with differing (but identically valued) base
    pointers would reference distinct regions of memory whereas the referenced
    memory regions were indeed identical.

    As part of this change we remove most of the MemoryAccess::get*BaseAddr
    interface. We removed already all references to get*BaseAddr in previous
    commits to ensure that no code relies on matching base pointers between
    memory accesses and scop arrays -- except for three remaining uses where we
    need the original base pointer. We document for these situations that
    MemoryAccess::getOriginalBaseAddr may return a base pointer that is distinct
    to the base pointer of the scop array referenced by this memory access.

Reviewers: sebpop, Meinersbur, zinob, gareevroman, pollydev, huihuiz, efriedma, jdoerfert

Reviewed By: Meinersbur

Subscribers: etherzhhb

Tags: #polly

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

llvm-svn: 302636
2017-05-10 10:59:58 +00:00
Siddharth Bhat a90be207c6 [Polly][PPCGCodeGen] OpenCL now gets kernel argument size from PPCG CodeGen
Summary: PPCGCodeGeneration now attaches the size of the kernel launch parameters at the end of the parameter list. For the existing CUDA Runtime, this gets ignored, but the OpenCL Runtime knows to check for kernel-argument size at the end of the parameter list. (The resulting parameters list is twice as long. This has been accounted for in the corresponding test cases).

Reviewers: grosser, Meinersbur, bollu

Reviewed By: bollu

Subscribers: nemanjai, yaxunl, Anastasia, pollydev, llvm-commits

Tags: #polly

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

llvm-svn: 302515
2017-05-09 10:45:52 +00:00
Siddharth Bhat 17f01968f1 [Polly] Added OpenCL Runtime to GPURuntime Library for GPGPU CodeGen
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
2017-05-07 21:03:46 +00:00
Tobias Grosser c6ad42165f Really disable test as intended in the previous commit
llvm-svn: 302360
2017-05-06 19:18:19 +00:00
Tobias Grosser 0f4e94673d Disable test to avoid buildbot noise
This test was introduced in r302339. It works on my system, but breaks on the
buildbots.

llvm-svn: 302358
2017-05-06 18:50:28 +00:00
Michael Kruse 5ae08c0ebb [DeLICM] Known knowledge.
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
2017-05-06 14:03:58 +00:00
Tobias Grosser c1ddedc657 Fix typo
llvm-svn: 302244
2017-05-05 15:46:01 +00:00
Michael Kruse f1052ceb5e [ScopBuilder] Do not verify unfeasible SCoPs.
SCoPs with unfeasible runtime context are thrown away and therefore
do not need their uses verified.

The added test case requires a complexity limit to exceed.
Normally, error statements are removed from the SCoP and for that
reason are skipped during the verification. If there is a unfeasible
runtime context (here: because of the complexity limit being reached),
the removal of error statements and other SCoP construction steps are
skipped to not waste time. Error statements are not modeled in SCoPs
and therefore have no requirements on whether the scalars used in
them are available.

llvm-svn: 302234
2017-05-05 13:38:35 +00:00