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
diff --git a/polly/lib/Analysis/ScopDetection.cpp b/polly/lib/Analysis/ScopDetection.cpp
index d1d6360..d5c154f 100644
--- a/polly/lib/Analysis/ScopDetection.cpp
+++ b/polly/lib/Analysis/ScopDetection.cpp
@@ -823,6 +823,15 @@
SmallVectorImpl<const SCEV *> &Sizes,
const SCEVUnknown *BasePointer,
Loop *Scope) const {
+ // If no sizes were found, all sizes are trivially valid. We allow this case
+ // to make it possible to pass known-affine accesses to the delinearization to
+ // try to recover some interesting multi-dimensional accesses, but to still
+ // allow the already known to be affine access in case the delinearization
+ // fails. In such situations, the delinearization will just return a Sizes
+ // array of size zero.
+ if (Sizes.size() == 0)
+ return true;
+
Value *BaseValue = BasePointer->getValue();
Region &CurRegion = Context.CurRegion;
for (const SCEV *DelinearizedSize : Sizes) {
@@ -893,10 +902,14 @@
else
IsNonAffine = true;
} else {
- SE.computeAccessFunctions(AF, Acc->DelinearizedSubscripts,
- Shape->DelinearizedSizes);
- if (Acc->DelinearizedSubscripts.size() == 0)
- IsNonAffine = true;
+ if (Shape->DelinearizedSizes.size() == 0) {
+ Acc->DelinearizedSubscripts.push_back(AF);
+ } else {
+ SE.computeAccessFunctions(AF, Acc->DelinearizedSubscripts,
+ Shape->DelinearizedSizes);
+ if (Acc->DelinearizedSubscripts.size() == 0)
+ IsNonAffine = true;
+ }
for (const SCEV *S : Acc->DelinearizedSubscripts)
if (!isAffine(S, Scope, Context))
IsNonAffine = true;
@@ -1013,7 +1026,7 @@
} else if (PollyDelinearize && !IsVariantInNonAffineLoop) {
Context.Accesses[BP].push_back({Inst, AF});
- if (!IsAffine)
+ if (!IsAffine || hasIVParams(AF))
Context.NonAffineAccesses.insert(
std::make_pair(BP, LI.getLoopFor(Inst->getParent())));
} else if (!AllowNonAffine && !IsAffine) {