If the access relation's domain is empty, the access will never be
executed. We can just remove it.
We only remove write accesses. Partial read accesses are not yet
supported and instructions in the statement might require the
llvm::Value holding the read's result to be defined.
llvm-svn: 308830
Change the indention of the last brace to align with the opening line.
Before:
Instructions {
%val = fadd double %arg, 2.100000e+01
store double %val, double* %A
}
After:
Instructions {
%val = fadd double %arg, 2.100000e+01
store double %val, double* %A
}
llvm-svn: 308828
Hoisted loads can be trivially supported because there are no
MemoryAccess to be modified, the loaded value is just available
at code generation.
llvm-svn: 308826
This pass 'forwards' operand trees into statements that use them in
order to avoid scalar dependencies.
This minimal implementation handles only the case of speculatable
instructions. We will successively add support for:
- Hoisted loads
- Read-only values
- Synthesizable values
- Loads
- PHIs
- Forwarding only parts of the tree
Differential Revision: https://reviews.llvm.org/D35754
llvm-svn: 308825
Summary:
For the ScopInfo lit testsuite, this patch removes some dependences on output behaviour of the legacy PM.
In most cases, these tests checked the tool output for labels created by the pass printer in the legacy PM. This doesn't work for the new PM anymore. Untangling the testcases is the first step to porting the testsuite for the new PM infrastructure.
Reviewers: grosser, Meinersbur, bollu
Reviewed By: grosser
Subscribers: llvm-commits, pollydev
Tags: #polly
Differential Revision: https://reviews.llvm.org/D35727
llvm-svn: 308754
Summary:
Added SPIR Code Generation to the PPCG Code Generator. This can be invoked using
the polly-gpu-arch flag value 'spir32' or 'spir64' for 32 and 64 bit code respectively.
In addition to that, runtime support has been added to execute said SPIR code on Intel
GPU's, where the system is equipped with Intel's open source driver Beignet (development
version). This requires the cmake flag 'USE_INTEL_OCL' to be turned on, and the polly-gpu-runtime
flag value to be 'libopencl'.
The transformation of LLVM IR to SPIR is currently quite a hack, consisting in part of regex
string transformations.
Has been tested (working) with Polybench 3.2 on an Intel i7-5500U (integrated graphics chip).
Reviewers: bollu, grosser, Meinersbur, singam-sanjay
Reviewed By: grosser, singam-sanjay
Subscribers: pollydev, nemanjai, mgorny, Anastasia, kbarton
Tags: #polly
Differential Revision: https://reviews.llvm.org/D35185
llvm-svn: 308751
Print a statement's instruction on dump() regardless of
-polly-print-instructions. dump() is supposed to be used in the debugger
only and never in regression tests. While debugging, get all the
information we have and we are not bound to break anything. For non-dump
purposes of print, forward the setting of -polly-print-instructions as
parameters.
Some calls to print() had to be changed because the
PollyPrintInstructions setting is only available in ScopInfo.cpp.
In ScheduleOptimizer.cpp, dump() was used in regression tests.
That's not what dump() is for.
The print parameter "PrintInstructions" will also be useful for an
explicit print SCoP pass in a future patch.
llvm-svn: 308746
When performing invariant load hoisting we check that invariant load expressions
are not too complex. Up to this commit, we performed this check by counting the
sum of dimensions in the access range as a very simple heuristic. This heuristic
is a little too conservative, as it prevents hoisting for any scops with a
very large number of parameters. Hence, we update the heuristic to only count
existentially quantified dimensions and set dimensions. We expect this to still
detect the problematic expressions in h264 because of which this check was
originally introduced.
For some unknown reason, this complexity check was originally committed in
IslNodeBuilder. It really belongs in ScopInfo, as there is no point in
optimizing a program which we could have known earlier cannot be code generated.
The benefit of running the check early is that we can avoid to even hoist checks
that are expensive to code generate as invariant loads. This can be seen in
the changed tests, where we now indeed detect the scop, but just not invariant
load hoist the complicated access.
We also improve the formatting of the code, document it, and use isl++ to
simplify expressions.
llvm-svn: 308659
Visual Studio, even the 2017 version, does not support C99 VLAs.
For VLA paramters, the length of the outermost dimension is not
required anyway, so remove it.
llvm-svn: 308643
When constructing a schedule true and there are multiple statements for
a basic block, create a sequence node for these statements.
Contributed-by: Nandini Singhal <cs15mtech01004@iith.ac.in>
Differential Revision: https://reviews.llvm.org/D35679
llvm-svn: 308635
We are working towards removing uses of Scop::getStmtFor(BB). In this
patch, we remove dependency of Scop::getLastStmtFor(BB) on
getStmtFor(BB). To do so, we get the list of all statements
corresponding to the BB and then fetch the last one.
Contributed-by: Nandini Singhal <cs15mtech01004@iith.ac.in>
Differential Revision: https://reviews.llvm.org/D35665
llvm-svn: 308633
Introduce previously missing PHIReads analogous the the already existing
PHIWrites/ValueWrites/ValueReads maps. PHIReads was initially not
required and the later introduced lookupPHIReadOf() used a linear
search instead.
With PHIReads, lookupPHIReadOf() can now also do a map lookup and remove
any surprising performance/behaviour differences to lookupPHIWriteOf(),
lookupValueWriteOf() and lookupValueReadOf().
llvm-svn: 308630
Use a mark-and-sweep algorithm to find and remove unused instructions
and MemoryAccesses. This is useful in particular to remove scalar
writes that are never used anywhere. A scalar write in a loop induces
a write-after-write dependency that stops the loop iterations to be
rescheduled. Such writes can be a result of previous transformations
such as DeLICM and operand tree forwarding.
It adds a new class VirtualInstruction that represents an instruction in
a particular statement. At the moment an instruction can only belong to
the statement that represents a BasicBlock. In the future, instructions
can be in one of multiple statements representing a BasicBlock
(Nandini's work), in different statements than its BasicBlock would
indicate, and even multiple statements at once (by forwarding operand
trees). It also integrates nicely with the VirtualUse class.
ScopStmt::contains(Instruction*) currently uses the instruction's parent
BasicBlock to check whether it contains the instruction. It will need to
check the actual statement list when one of the aforementioned features
become possible.
Differential Revision: https://reviews.llvm.org/D35656
llvm-svn: 308626
This commit *WILL COMPILE*.
1. `PPCG` now uses `isl_multi_pw_aff` instead of an array of `pw_aff`.
This needs us to adjust how we index array bounds and how we construct
array bounds.
2. `PPCG` introduces two new kinds of nodes: `init_device` and `clear_device`.
We should investigate what the correct way to handle these are.
3. `PPCG` has gotten smarter with its use of live range reordering, so some of
the tests have a qualitative improvement.
4. `PPCG` changed its output style, so many test cases need to be updated to
fit the new style for `polly-acc-dump-code` checks.
Differential Revision: https://reviews.llvm.org/D35677
llvm-svn: 308625
- This commit *WILL NOT COMPILE*. `PPCGCodeGeneration` requires changes
since some of PPCG's internal data structures have been modified.
- Has polly-speific changes to PPCG. Polly exports certain functionality that
is private to PPCG. It also creates stubs for large parts of the pet API as
well as other functions in `ppcg/external.c` to keep the linker happy.
- This commit includes changes to CMakeLists.txt.
Differential Revision: https://reviews.llvm.org/D35676
llvm-svn: 308624
- This commit *WILL NOT COMPILE*, as it checks in vanilla PPCG 0.07
- We choose to introduce this commit into the history to cleanly display
the Polly-specific changes made to PPCG.
Differential Revision: https://reviews.llvm.org/D35675
llvm-svn: 308623
Since there will be no more a 1:1 correspondence between statements
and basic blocks, we would like to get rid of the method getStmtFor(BB)
and its uses. Here we remove one of its uses in ScopBuilder by fetching
the statement in which the instruction lies.
Contributed-by: Nandini Singhal <cs15mtech01004@iith.ac.in>
Differential Revision: https://reviews.llvm.org/D35610
llvm-svn: 308610
This is one possible solution to implement wrap-arounds for integers in
unsigned icmp operations. For example,
store i32 -1, i32* %A_addr
%0 = load i32, i32* %A_addr
%1 = icmp ult i32 %0, 0
%1 should hold false, because under the assumption of unsigned integers,
-1 should wrap around to 2^32-1. However, previously. it was assumed
that the MSB (Most Significant Bit - aka the Sign bit) was never set for
integers in unsigned operations.
This patch modifies the buildConditionSets function in ScopInfo.cpp to
give better information about the integers in these unsigned
comparisons.
Contributed-by: Annanay Agarwal <cs14btech11001@iith.ac.in>
Differential Revision: https://reviews.llvm.org/D35464
llvm-svn: 308608
Before this patch, ScalarDefUseChain was a tool used by DeLICM to find
all reads and writes of scalar accesses. It iterated once over all
accesses and stores the accesses into maps.
By integrating it into the Scop class, we can keep the maps up-to-date
without the need for recomputing them. It will be needed for more than
DeLICM in the future, such as SCoP simplification, code movement between
virtual statements, and array expansion (GSoC project).
Compared to ScalarUseDefChain, we save two maps by finding the ScopStmt
a Def/PHIRead must reside in, and use its already existing lookup
function to find the MemoryAccess.
Differential Revision: https://reviews.llvm.org/D35631
llvm-svn: 308495
Some optimizations (e.g., DeLICM) can modify memory accesses (e.g., change
their MemoryKind). Consequently, the pattern matching should take it into
the account.
Reviewed-by: Tobias Grosser <tobias@grosser.es>,
Michael Kruse <llvm@meinersbur.de>
Differential Revision: https://reviews.llvm.org/D33138
llvm-svn: 308494
Summary:
This makes code more readable and allows to reuse this functionality in
the future at other places.
Suggested-by Michael Kruse in post-commit review of r307660.
Reviewers: Meinersbur, bollu, gareevroman, efriedma, huihuiz, sebpop, simbuerg
Reviewed By: Meinersbur
Subscribers: pollydev, llvm-commits
Tags: #polly
Differential Revision: https://reviews.llvm.org/D35585
llvm-svn: 308435
Once statements are split, a BasicBlock will comprise of multiple
statements. To prepare for this change in future, we introduce a list
of statements in the statement map.
Contributed-by: Nandini Singhal <cs15mtech01004@iith.ac.in>
Differential Revision: https://reviews.llvm.org/D35301
llvm-svn: 308318
find_library(lib) stores the result in the variable "lib", which is also
cached in CMakeCache.txt. This could theoretically conflict if jsoncpp
required two libraries, which each would get cached as "lib". Use a
more descriptive and disambiguative "jsoncpp_${libname}" for that.
llvm-svn: 308289
Use ${libname} instead of ${lib}. By a coincidence, this worked
because ${lib} also the variable used for finding the libjsoncpp.so
full path.
llvm-svn: 308288
pkg_search_module(JSONCPP) should set
JSONCPP_LIBDIR/JSONCPP_LIBRARY_DIRS to where the libjsoncpp.so can be
found. However, on Ubuntu 14.04 LTS (Trusty Tahr) it returns /usr/lib
while the libjsoncpp library can be found at
/usr/lib/x86_64-linux-gnu/libjsoncpp.so. Thus, while searching for
the full path of the jsoncpp library, it is not found.
JSONCPP_LIBDIR is correctly set to /usr/lib/x86_64-linux-gnu on e.g.,
Ubuntu 16.04 LTS (Xenial Xerus )
Fix by removing the NO_DEFAULT_PATH flag, in order to search the system
default paths even if the library is not found in
JSONCPP_LIBDIR/JSONCPP_LIBRARY_DIRS.
This fixes bug llvm.org/PR33798.
llvm-svn: 308287
We extended kills in Polly to handle both `phi` nodes and scalars that
are not used within the Scop. Update the comments and choice of
variable names to reflect this.
llvm-svn: 308279
Utilizing newer LLVM diagnostic remark API in order to enable use of
opt-viewer tool. Polly Diagnostic Remarks also now appear in YAML
remark file.
In this patch, I've added the OptimizationRemarkEmitter into certain
classes where remarks are being emitted and update the remark emit calls
itself. I also provide each remark a BasicBlock or Instruction from where
it is being called, in order to compute the hotness of the remark.
Patch by Tarun Rajendran!
Differential Revision: https://reviews.llvm.org/D35399
llvm-svn: 308233
Summary: Since there will be no more a 1-1 correspondence between statements and basic block, we would like to get rid of the method `getStmtFor(BB)` and its uses. Here we remove one of its uses in PolyhedralInfo, as suggested by Michael Sir.
Reviewers: grosser, Meinersbur, bollu
Reviewed By: grosser
Subscribers: pollydev
Tags: #polly
Differential Revision: https://reviews.llvm.org/D35300
llvm-svn: 308220
Summary:
We do not keep domain constraints on access functions when building the
scop. Hence, for consistency reasons, it makes also sense to not include
them when storing a new access function. This change results in simpler
access functions that make output easier to read.
This patch also helps to make DeLICMed memory accesses to be understood by
our matrix multiplication pattern matching pass. Further changes to the
matrix multiplication pattern matching are needed for this to work, so the
corresponding test case will be added in a future commit.
Reviewers: Meinersbur, bollu, gareevroman, efriedma, huihuiz, sebpop, simbuerg
Subscribers: pollydev, llvm-commits
Tags: #polly
Differential Revision: https://reviews.llvm.org/D35237
llvm-svn: 308215
- We should call `preloadInvariantLoads` to make sure that code is
generated for invariant loads in the kernel.
Differential Revision: https://reviews.llvm.org/D35410
llvm-svn: 308187
This patch makes sure that in case a loop is not fully contained within a region
that later forms a SCoP, none of the loop backedges are allowed to be part of
the region. We currently do not support the situation where only some of a loops
backedges are part of a scop. Today, this can break both scop modeling and code
generation. One such breaking test case is for example
test/ScopDetectionDiagnostics/loop_partially_in_scop-2.ll, where we totally
forgot to code generate some of the backedges. Fortunately, it is commonly not
necessary to support these partial loops, it is way more common that either
no backedge is included in a region or all loop backedge are included.
This fixes a recent miscompile in
MultiSource/Benchmarks/MiBench/consumer-typeset which was exposed after
r306477.
llvm-svn: 308113
- There is a conditional branch that is used to switch between the old
and new versions of the code.
- If we detect that the build was unsuccessful, `PPCGCodeGeneration` will
change the runtime check to be always set to false.
- To actually *reach* this runtime check instruction, `PPCGCodeGeneration`
was using assumptions about the layout of the BBs.
- However, invariant load hoisting violates this assumption by inserting
an extra basic block in the middle.
- Fix the assumption on the layout by having `createScopConditionally`
return the conditional branch instruction.
- Use this reference to set to always-false.
llvm-svn: 308010
We need to relax constraints on invariant loads so that they do not
create fake RAW dependences. So, we do not consider invariant loads as
scalar dependences in a region.
During these changes, it turned out that we do not consider `llvm::Value`
replacements correctly within `PPCGCodeGeneration` and `ISLNodeBuilder`.
The replacements dictated by `ValueMap` were not being followed in all
places. This was fixed in this commit. There is no clean way to decouple
this change because this bug only seems to arise when the relaxed
version of invariant load hoisting was enabled.
Differential Revision: https://reviews.llvm.org/D35120
llvm-svn: 307907
Summary:
Add a sequence number that identifies a ptx_kernel's parent Scop within a function to it's name to differentiate it from other kernels produced from the same function, yet different Scops.
Kernels produced from different Scops can end up having the same name. Consider a function with 2 Scops and each Scop being able to produce just one kernel. Both of these kernels have the name "kernel_0". This can lead to the wrong kernel being launched when the runtime picks a kernel from its cache based on the name alone. This patch supplements D33985, by differentiating kernels across Scops as well.
Previously (even before D33985) while profiling kernels generated through JIT e.g. Julia, [[ https://groups.google.com/d/msg/polly-dev/J1j587H3-Qw/mR-jfL16BgAJ | kernels associated with different functions, and even different SCoPs within a function, would be grouped together due to the common name ]]. This patch prevents this grouping and the kernels are reported separately.
Reviewers: grosser, bollu
Reviewed By: grosser
Subscribers: mehdi_amini, nemanjai, pollydev, kbarton
Tags: #polly
Differential Revision: https://reviews.llvm.org/D35176
llvm-svn: 307814
- `lit.util.capture` was removed in `r306625`.
- Replace `lit.util.capture` to `subprocess.check_output` as LLVM did.
- LLVM revision of this change: `https://reviews.llvm.org/D35088`.
Differential Revision: https://reviews.llvm.org/D35255
llvm-svn: 307765
Summary:
There is a bug in the current lit configurations for the unittests. If gtest is not available, the site-config for the unit tests won't be generated. Because lit recurses through the test directory, the lit configuration for the unit tests will be discovered nevertheless, leading to a fatal error in lit.
This patch semi-gracefully skips the unittests if gtest is not available. As a result, running lit now prints this: `warning: test suite 'Polly-Unit' contained no test`.
If people think that this is too annoying, the alternative would be to pick apart the test directory, so that the lit testsuite discovery will always only find one configuration. In fact, both of these things could be combined. While it's certainly nice that running a single lit command runs all the tests, I suppose people use the `check-polly` make target over lit most of the time, so the difference might not be noticed.
Reviewers: Meinersbur, grosser
Reviewed By: grosser
Subscribers: mgorny, bollu, pollydev, llvm-commits
Tags: #polly
Differential Revision: https://reviews.llvm.org/D34053
llvm-svn: 307651
Summary:
As of now, Polly uses llvm-config to set up LLVM dependencies in an out-of-tree build.
This is problematic for two reasons:
1) Right now, in-tree and out-of-tree builds in fact do different things. E.g., in an in-tree build, libPolly depends on a handful of LLVM libraries, while in an out-of-tree build it depends on all of them. This means that we often need to treat both paths seperately.
2) I'm specifically unhappy with the way libPolly is linked right now, because it just blindly links against all the LLVM libs. That doesn't make a lot of sense. For instance, one of these libs is LLVMTableGen, which contains a command line definition of a -o option. This means that I can not link an out-of-tree libPolly into a tool which might want to offer a -o option as well.
This patch (mostly) drop the use of llvm-config in favor of LLVMs exported cmake package. However, building Polly with unittests requires access to the gtest sources (in the LLVM source tree). If we're building against an LLVM installation, this source tree is unavailable and must specified. I'm using llvm-config to provide a default in this case.
Reviewers: Meinersbur, grosser
Reviewed By: grosser
Subscribers: tstellar, bollu, chapuni, mgorny, pollydev, llvm-commits
Differential Revision: https://reviews.llvm.org/D33299
llvm-svn: 307650
For the previous commit I accidentally added this change to lit.site.cfg, which
is autogenerated and was consequently not part of the previous commit.
llvm-svn: 307648
When providing the option "-polly-ast-print-accesses" Polly also prints the
memory accesses that are generated:
#pragma known-parallel
for (int c0 = 0; c0 <= 1023; c0 += 4)
#pragma simd
for (int c1 = c0; c1 <= c0 + 3; c1 += 1)
Stmt_for_body(
/* read */ &MemRef_B[0]
/* write */ MemRef_A[c1]
);
This makes writing and debugging memory layout transformations easier.
Based on a patch contributed by Thomas Lang (ETH Zurich)
llvm-svn: 307579
Summary:
Since r306667, propagateInvalidStmtDomains gets a reference to an
InvalidDomainMap. As part of the branch leading to return false, the respective
domain is freed. It is, however, not removed from the InvalidDomainMap, leaking
a pointer to a freed object which results in a use-after-free. Fix this be
removing the domain from the map before returning.
We tried to derive a test case that reliably failes, but did not succeed in
producing one. Hence, for now the failures in our LNT bots must be sufficient
to keep this issue tested.
Reviewers: grosser, Meinersbur, bollu
Subscribers: bollu, nandini12396, pollydev, llvm-commits
Differential Revision: https://reviews.llvm.org/D34971
llvm-svn: 307499
- Check that we have invariant accesses.
- Use `-polly-use-llvm-names` for better names in the test.
- Rename test function to `f` for brevity.
llvm-svn: 307401
- This already works, but add this to ensure that there is no
regressions when I expand the invariant load hoisting ability of
`PPCGCodeGeneration`.
llvm-svn: 307398
- Instead of running with -O0, we enable the highest optimization level, but
then disable optimizations. This ensures that possibly important metadata
is still emitted.
- Update the code for attribute removal to work with latest LLVM
- Do not cut an arbitrary number of lines from the LL file. It is undocumented
why this was needed at the first place, and such a feature is likely to
break with trivial IR changes that may come in the future.
llvm-svn: 307355
- By definition, we can pass something as a `kill` to PPCG if we know
that no data can flow across a kill.
- This is useful for more complex examples where we have scalars that
are local to a scop.
- If the local is only used within a scop, we are free to kill it.
Differential Revision: https://reviews.llvm.org/D35045
llvm-svn: 307260
Summary:
Provide more context to the name of a GPU kernel by prefixing its name with the host function that calls it. E.g. The first kernel called by `gemm` would be `FUNC_gemm_KERNEL_0`.
Kernels currently follow the "kernel_#" (# = 0,1,2,3,...) nomenclature. This patch makes it easier to map host caller and device callee, especially when there are many kernels produced by Polly-ACC.
Reviewers: grosser, Meinersbur, bollu, philip.pfaffe, kbarton!
Reviewed By: grosser
Subscribers: nemanjai, pollydev
Tags: #polly
Differential Revision: https://reviews.llvm.org/D33985
llvm-svn: 307173
Polly did not use PPCG's live range reordering feature. Teach
PPCGCodeGeneration to use this.
Documentation on this is sparse, so much of the code is conservative.
We currently kill all phi nodes in a Scop by appending them to the
must_kill map we pass to PPCG. I do not have a proof of correctness,
but it seems to be intuitively correct.
We also do not handle `array_order`, which, quoting PPCG, is:
PPCG/gpu.h: "Order dependences on non-scalars."
It seems to consist of RAW dependences between arrays. We need to
pass this information for more complex privatization cases.
Differential Revision: https://reviews.llvm.org/D34941
llvm-svn: 307163
Summary: This is a general maintenance update
Reviewers: grosser
Subscribers: srhines, fedor.sergeev, pollydev, llvm-commits
Contributed-by: Maximilian Falkenstein <falkensm@student.ethz.ch>
Differential Revision: https://reviews.llvm.org/D34903
llvm-svn: 307090
Summary:
Introduce a "hybrid" `-polly-target` option to optimise code for either the GPU or CPU.
When this target is selected, PPCGCodeGeneration will attempt first to optimise a Scop. If the Scop isn't modified, it is then sent to the passes that form the CPU pipeline, i.e. IslScheduleOptimizerPass, IslAstInfoWrapperPass and CodeGeneration.
In case the Scop is modified, it is marked to be skipped by the subsequent CPU optimisation passes.
Reviewers: grosser, Meinersbur, bollu
Reviewed By: grosser
Subscribers: kbarton, nemanjai, pollydev
Tags: #polly
Differential Revision: https://reviews.llvm.org/D34054
llvm-svn: 306863
this is a great test file name based on this update, but I'll let Polly
folks sort out how they want this to work long-term, I just want tho
bots back.
llvm-svn: 306767
ScopStmts were being used in the computation of the Domain of the SCoPs
in ScopInfo. Once statements are split, there will not be a 1-to-1
correspondence between Stmts and Basic blocks. Thus this patch avoids
the use of getStmtFor() by creating a map of BB to InvalidDomain and
using it to compute the domain of the statements.
Contributed-by: Nanidini Singhal <cs15mtech01004@iith.ac.in>
Differential Revision: https://reviews.llvm.org/D33942
llvm-svn: 306667
Summary:
The NVPTX backend is now initialised within Polly. A language front-end need not be modified to initialise the backend, just for Polly.
Reviewers: Meinersbur, grosser
Reviewed By: Meinersbur
Subscribers: vchuravy, mgorny
Tags: #polly
Differential Revision: https://reviews.llvm.org/D31859
llvm-svn: 306649
This patch aims to implement the option of allocating new arrays created
by polly on heap instead of stack. To enable this option, a key named
'allocation' must be written in the imported json file with the value
'heap'.
We need such a feature because in a next iteration, we will implement a
mechanism of maximal static expansion which will need a way to allocate
arrays on heap. Indeed, the expansion is very costly in terms of memory
and doing the allocation on stack is not worth considering.
The malloc and the free are added respectively at polly.start and
polly.exiting such that there is no use-after-free (for instance in case
of Scop in a loop) and such that all memory cells allocated with a
malloc are free'd when we don't need them anymore.
We also add :
- In the class ScopArrayInfo, we add a boolean as member called IsOnHeap
which represents the fact that the array in allocated on heap or not.
- A new branch in the method allocateNewArrays in the ISLNodeBuilder for
the case of heap allocation. allocateNewArrays now takes a BBPair
containing polly.start and polly.exiting. allocateNewArrays takes this
two blocks and add the malloc and free calls respectively to
polly.start and polly.exiting.
- As IntPtrTy for the malloc call, we use the DataLayout one.
To do that, we have modified :
- createScopArrayInfo and getOrCreateScopArrayInfo such that it returns
a non-const SAI, in order to be able to call setIsOnHeap in the
JSONImporter.
- executeScopConditionnaly such that it return both start block and end
block of the scop, because we need this two blocs to be able to add
the malloc and the free calls at the right position.
Differential Revision: https://reviews.llvm.org/D33688
llvm-svn: 306540
Before we would 'guess' the correct location for the MergeBlock
that got introduced when executing a Scop conditionally. This
implicitly depends on the situation that at this point during
CodeGen there will be nothing between polly.start and polly.exiting.
With this commit we explicitly state that we want the block that
directly follows polly.exiting.
llvm-svn: 306398
This test fails, if polly is not linked into LLVM's tools. Our
lit site-config already deals with this by not adding the -load
option, if polly is linked into LLVM's tools.
llvm-svn: 306395
- In D33414, if any function call was found within a kernel, we would bail out.
- This is an over-approximation. This patch changes this by allowing the
`llvm.sqrt.*` family of intrinsics.
- This introduces an additional step when creating a separate llvm::Module
for a kernel (GPUModule). We now copy function declarations from the
original module to new module.
- We also populate IslNodeBuilder::ValueMap so it replaces the function
references to the old module to the ones in the new module
(GPUModule).
Differential Revision: https://reviews.llvm.org/D34145
llvm-svn: 306284
This commit returns both the start and the exit block that are created
by executeScopConditionally.
In a future commit we will make use of the exit block. Before we would
have to use the implicit property that there won't be any code generated
between polly.start and polly.exiting at the time of use to find the
correct block ('polly.exiting').
All usage location are semantically unchanged.
llvm-svn: 306283
The condition that disallowed code generation in PPCGCodeGeneration with
invariant loads is not required. I haven't been able to construct a
counterexample where this generates invalid code.
Differential Revision: https://reviews.llvm.org/D34604
llvm-svn: 306245
This reduces the compilation time of one reduced test case from Android from
16 seconds to 100 mseconds (we bail out), without negatively impacting any
other test case we currently have.
We still saw occasionally compilation timeouts on the AOSP buildbot. Hopefully,
those will go away with this change.
llvm-svn: 306235
During the construction of MemoryAccesses in ScopBuilder, BasicBlocks
were used in function parameters, assuming that the ScopStmt an be
directly derived from it. This won't be true anymore once we split
BasicBlocks into multiple ScopStmt. As a preparation for such a change
in the future, we instead pass the ScopStmt and avoid the use of
getStmtFor().
There are two occasions where a kind of mapping from BasicBlock to
ScopStmt is still required.
1. Get the statement representing the incoming block of a `PHINode`
using `getLastStmtOf`.
2. One statement is required to write a scalar to be readable by those
which need it. This is most often the statement which contains its
definition, which we get using `getStmtFor(Instruction*)`.
Differential Revision: https://reviews.llvm.org/D34369
llvm-svn: 306132
This allows us to bail out both in case the lexmin/max computation is too
expensive, but also in case the commulative cost across an alias group is
too expensive. This is an improvement of r303404, which did not seem to
be sufficient to keep the Android Buildbot quiet.
llvm-svn: 306087
r303971 added an assertion that SCEV addition involving an AddRec
and a SCEVUnknown must involve a dominance relation: either the
SCEVUnknown value dominates the AddRec's loop, or the AddRec's
loop header dominates the SCEVUnknown. This is generally fine
for most usage of SCEV because it isn't possible to write an
expression in IR which would violate it, but it's a bit inconvenient
here for polly.
To solve the issue, just avoid creating a SCEV expression which
triggers the asssertion.
I'm not really happy with this solution, but I don't have any better
ideas.
Fixes https://bugs.llvm.org/show_bug.cgi?id=33464.
Differential Revision: https://reviews.llvm.org/D34259
llvm-svn: 305864
Ensure that all array base pointers are assigned before generating
aliasing metadata by allocating new arrays beforehand.
Before this patch, getBasePtr() returned nullptr for new arrays because
the arrays were created at a later point. Nullptr did not match to any
array after the created array base pointers have been assigned and when
the loads/stores are generated.
llvm-svn: 305675
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
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
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
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
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
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
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
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
- 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
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
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
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
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
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
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
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
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
- 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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
This commit exports the majority of the isl functions to the isl C++ interface.
The official isl C++ bindings still require discussions to define the set of
functions that are officially supported. As a result, the officially exported
functionality will be rather limited until these discussions conclude and a
non-trivial set of isl functions is officially supported through the isl C++
bindings. Starting from this commit we ship with Polly an extended version of
the official isl C++ bindings to ensure sufficient functionality is available
such that LLVM developers can make efficient use of isl through C++. The
practical experience Polly gathers with its bindings will then be used to
gradually upstream patches to isl to extend the official bindings.
llvm-svn: 303506
This reduces the diff to the official isl C++ bindings and solves a correctness
issue with isl::booleans, where isl_bool_error results were accidentally
converted to isl::boolean::true.
llvm-svn: 303505
Instead of relying on these functions to be part of the isl C++ bindings, we
just define this functionality independently. This allows us to use isl C++
bindings that do not contain LLVM specific functionality.
llvm-svn: 303503
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
- 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
In r302231 we mistakenly use bitwise or (|) instead of logical
or (||). This patch fixes that.
Contributed-by: Sameer AbuAsal <sabuasal@codeaurora.org>
Differential Revision: https://reviews.llvm.org/D33337
llvm-svn: 303386
Summary:
Implements PR889
Removing the virtual table pointer from Value saves 1% of RSS when doing
LTO of llc on Linux. The impact on time was positive, but too noisy to
conclusively say that performance improved. Here is a link to the
spreadsheet with the original data:
https://docs.google.com/spreadsheets/d/1F4FHir0qYnV0MEp2sYYp_BuvnJgWlWPhWOwZ6LbW7W4/edit?usp=sharing
This change makes it invalid to directly delete a Value, User, or
Instruction pointer. Instead, such code can be rewritten to a null check
and a call Value::deleteValue(). Value objects tend to have their
lifetimes managed through iplist, so for the most part, this isn't a big
deal. However, there are some places where LLVM deletes values, and
those places had to be migrated to deleteValue. I have also created
llvm::unique_value, which has a custom deleter, so it can be used in
place of std::unique_ptr<Value>.
I had to add the "DerivedUser" Deleter escape hatch for MemorySSA, which
derives from User outside of lib/IR. Code in IR cannot include MemorySSA
headers or call the MemoryAccess object destructors without introducing
a circular dependency, so we need some level of indirection.
Unfortunately, no class derived from User may have any virtual methods,
because adding a virtual method would break User::getHungOffOperands(),
which assumes that it can find the use list immediately prior to the
User object. I've added a static_assert to the appropriate OperandTraits
templates to help people avoid this trap.
Reviewers: chandlerc, mehdi_amini, pete, dberlin, george.burgess.iv
Reviewed By: chandlerc
Subscribers: krytarowski, eraman, george.burgess.iv, mzolotukhin, Prazek, nlewycky, hans, inglorion, pcc, tejohnson, dberlin, llvm-commits
Differential Revision: https://reviews.llvm.org/D31261
llvm-svn: 303362
Michael Kruse