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

1 files changed, 18 insertions, 5 deletions

diff --git a/polly/lib/Analysis/ScopDetection.cpp b/polly/lib/Analysis/ScopDetection.cpp
index d1d6360ab44..d5c154fa08a 100644
--- a/polly/lib/Analysis/ScopDetection.cpp
+++ b/polly/lib/Analysis/ScopDetection.cpp
@@ -823,6 +823,15 @@ bool ScopDetection::hasValidArraySizes(DetectionContext &Context,
                                        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 @@ bool ScopDetection::computeAccessFunctions(
       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 @@ bool ScopDetection::isValidAccess(Instruction *Inst, const SCEV *AF,
   } 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) {