summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
Diffstat
-rw-r--r--polly/include/polly/ScopDetection.h1
-rw-r--r--polly/include/polly/TempScopInfo.h10
-rw-r--r--polly/lib/Analysis/ScopDetection.cpp27
-rw-r--r--polly/lib/Analysis/ScopInfo.cpp42
-rw-r--r--polly/lib/Analysis/TempScopInfo.cpp49
-rw-r--r--polly/test/Dependences/do_pluto_matmult.ll2
-rw-r--r--polly/test/Isl/Ast/OpenMP/nested_loop_both_parallel_parametric.ll4
-rw-r--r--polly/test/ScopInfo/constant_start_integer.ll31
-rw-r--r--polly/test/ScopInfo/multidim_ivs_and_integer_offsets_3d.ll35
-rw-r--r--polly/test/ScopInfo/multidim_ivs_and_parameteric_offsets_3d.ll30
-rw-r--r--polly/test/ScopInfo/multidim_nested_start_integer.ll31
-rw-r--r--polly/test/ScopInfo/multidim_nested_start_share_parameter.ll23
-rw-r--r--polly/test/ScopInfo/multidim_only_ivs_2d.ll30
-rw-r--r--polly/test/ScopInfo/multidim_only_ivs_3d.ll30
-rw-r--r--polly/test/ScopInfo/multidim_only_ivs_3d_cast.ll33
15 files changed, 202 insertions, 176 deletions
diff --git a/polly/include/polly/ScopDetection.h b/polly/include/polly/ScopDetection.h
index defc2f175df..452eb09e49b 100644
--- a/polly/include/polly/ScopDetection.h
+++ b/polly/include/polly/ScopDetection.h
@@ -73,6 +73,7 @@ namespace polly {
typedef std::set<const SCEV *> ParamSetType;
extern bool PollyTrackFailures;
+extern bool PollyDelinearize;
//===----------------------------------------------------------------------===//
/// @brief Pass to detect the maximal static control parts (Scops) of a
diff --git a/polly/include/polly/TempScopInfo.h b/polly/include/polly/TempScopInfo.h
index 5119157148f..7e447f55099 100644
--- a/polly/include/polly/TempScopInfo.h
+++ b/polly/include/polly/TempScopInfo.h
@@ -29,6 +29,8 @@ using namespace llvm;
namespace polly {
+extern bool PollyDelinearize;
+
//===---------------------------------------------------------------------===//
/// @brief A memory access described by a SCEV expression and the access type.
class IRAccess {
@@ -52,10 +54,14 @@ private:
bool IsAffine;
public:
+ SmallVector<const SCEV *, 4> Subscripts, Sizes;
+
explicit IRAccess(TypeKind Type, const Value *BaseAddress, const SCEV *Offset,
- unsigned elemBytes, bool Affine)
+ unsigned elemBytes, bool Affine,
+ SmallVector<const SCEV *, 4> Subscripts,
+ SmallVector<const SCEV *, 4> Sizes)
: BaseAddress(BaseAddress), Offset(Offset), ElemBytes(elemBytes),
- Type(Type), IsAffine(Affine) {}
+ Type(Type), IsAffine(Affine), Subscripts(Subscripts), Sizes(Sizes) {}
enum TypeKind getType() const { return Type; }
diff --git a/polly/lib/Analysis/ScopDetection.cpp b/polly/lib/Analysis/ScopDetection.cpp
index 69c63ae1b13..5f03e7a5b45 100644
--- a/polly/lib/Analysis/ScopDetection.cpp
+++ b/polly/lib/Analysis/ScopDetection.cpp
@@ -117,6 +117,12 @@ TrackFailures("polly-detect-track-failures",
cl::location(PollyTrackFailures), cl::Hidden, cl::ZeroOrMore,
cl::init(false), cl::cat(PollyCategory));
+static cl::opt<bool, true>
+PollyDelinearizeX("polly-delinearize",
+ cl::desc("Delinearize array access functions"),
+ cl::location(PollyDelinearize), cl::Hidden, cl::ZeroOrMore,
+ cl::init(false), cl::cat(PollyCategory));
+
static cl::opt<bool>
VerifyScops("polly-detect-verify",
cl::desc("Verify the detected SCoPs after each transformation"),
@@ -124,6 +130,7 @@ VerifyScops("polly-detect-verify",
cl::cat(PollyCategory));
bool polly::PollyTrackFailures = false;
+bool polly::PollyDelinearize = false;
//===----------------------------------------------------------------------===//
// Statistics.
@@ -357,11 +364,25 @@ bool ScopDetection::isValidMemoryAccess(Instruction &Inst,
return invalid<ReportVariantBasePtr>(Context, /*Assert=*/false, BaseValue);
AccessFunction = SE->getMinusSCEV(AccessFunction, BasePointer);
-
- if (!AllowNonAffine &&
- !isAffineExpr(&Context.CurRegion, AccessFunction, *SE, BaseValue))
+ const SCEVAddRecExpr *AF = dyn_cast<SCEVAddRecExpr>(AccessFunction);
+
+ if (AllowNonAffine) {
+ // Do not check whether AccessFunction is affine.
+ } else if (PollyDelinearize && AF) {
+ // Try to delinearize AccessFunction.
+ SmallVector<const SCEV *, 4> Subscripts, Sizes;
+ AF->delinearize(*SE, Subscripts, Sizes);
+ int size = Subscripts.size();
+
+ for (int i = 0; i < size; ++i)
+ if (!isAffineExpr(&Context.CurRegion, Subscripts[i], *SE, BaseValue))
+ return invalid<ReportNonAffineAccess>(Context, /*Assert=*/true,
+ AccessFunction);
+ } else if (!isAffineExpr(&Context.CurRegion, AccessFunction, *SE,
+ BaseValue)) {
return invalid<ReportNonAffineAccess>(Context, /*Assert=*/true,
AccessFunction);
+ }
// FIXME: Alias Analysis thinks IntToPtrInst aliases with alloca instructions
// created by IndependentBlocks Pass.
diff --git a/polly/lib/Analysis/ScopInfo.cpp b/polly/lib/Analysis/ScopInfo.cpp
index fe8bafe8b91..245d68311ef 100644
--- a/polly/lib/Analysis/ScopInfo.cpp
+++ b/polly/lib/Analysis/ScopInfo.cpp
@@ -337,21 +337,37 @@ MemoryAccess::MemoryAccess(const IRAccess &Access, const Instruction *AccInst,
Type = Access.isRead() ? READ : MUST_WRITE;
- isl_pw_aff *Affine = SCEVAffinator::getPwAff(Statement, Access.getOffset());
+ int Size = Access.Subscripts.size();
+ assert(Size > 0 && "access function with no subscripts");
+ AccessRelation = NULL;
+
+ for (int i = 0; i < Size; ++i) {
+ isl_pw_aff *Affine =
+ SCEVAffinator::getPwAff(Statement, Access.Subscripts[i]);
+
+ if (i == Size - 1) {
+ // Divide the access function of the last subscript by the size of the
+ // elements in the array.
+ //
+ // A stride one array access in C expressed as A[i] is expressed in
+ // LLVM-IR as something like A[i * elementsize]. This hides the fact that
+ // two subsequent values of 'i' index two values that are stored next to
+ // each other in memory. By this division we make this characteristic
+ // obvious again.
+ isl_val *v;
+ v = isl_val_int_from_si(isl_pw_aff_get_ctx(Affine),
+ Access.getElemSizeInBytes());
+ Affine = isl_pw_aff_scale_down_val(Affine, v);
+ }
- // Divide the access function by the size of the elements in the array.
- //
- // A stride one array access in C expressed as A[i] is expressed in LLVM-IR
- // as something like A[i * elementsize]. This hides the fact that two
- // subsequent values of 'i' index two values that are stored next to each
- // other in memory. By this division we make this characteristic obvious
- // again.
- isl_val *v;
- v = isl_val_int_from_si(isl_pw_aff_get_ctx(Affine),
- Access.getElemSizeInBytes());
- Affine = isl_pw_aff_scale_down_val(Affine, v);
+ isl_map *SubscriptMap = isl_map_from_pw_aff(Affine);
+
+ if (!AccessRelation)
+ AccessRelation = SubscriptMap;
+ else
+ AccessRelation = isl_map_flat_range_product(AccessRelation, SubscriptMap);
+ }
- AccessRelation = isl_map_from_pw_aff(Affine);
isl_space *Space = Statement->getDomainSpace();
AccessRelation = isl_map_set_tuple_id(
AccessRelation, isl_dim_in, isl_space_get_tuple_id(Space, isl_dim_set));
diff --git a/polly/lib/Analysis/TempScopInfo.cpp b/polly/lib/Analysis/TempScopInfo.cpp
index 7eff1260ecb..62dc47cd69f 100644
--- a/polly/lib/Analysis/TempScopInfo.cpp
+++ b/polly/lib/Analysis/TempScopInfo.cpp
@@ -131,9 +131,14 @@ bool TempScopInfo::buildScalarDependences(Instruction *Inst, Region *R) {
assert(!isa<PHINode>(UI) && "Non synthesizable PHINode found in a SCoP!");
+ SmallVector<const SCEV *, 4> Subscripts, Sizes;
+ Subscripts.push_back(SE->getConstant(ZeroOffset->getType(), 0));
+ Sizes.push_back(SE->getConstant(ZeroOffset->getType(), 1));
+
// Use the def instruction as base address of the IRAccess, so that it will
// become the name of the scalar access in the polyhedral form.
- IRAccess ScalarAccess(IRAccess::SCALARREAD, Inst, ZeroOffset, 1, true);
+ IRAccess ScalarAccess(IRAccess::SCALARREAD, Inst, ZeroOffset, 1, true,
+ Subscripts, Sizes);
AccFuncMap[UseParent].push_back(std::make_pair(ScalarAccess, UI));
}
@@ -142,14 +147,17 @@ bool TempScopInfo::buildScalarDependences(Instruction *Inst, Region *R) {
IRAccess TempScopInfo::buildIRAccess(Instruction *Inst, Loop *L, Region *R) {
unsigned Size;
+ Type *SizeType;
enum IRAccess::TypeKind Type;
if (LoadInst *Load = dyn_cast<LoadInst>(Inst)) {
- Size = TD->getTypeStoreSize(Load->getType());
+ SizeType = Load->getType();
+ Size = TD->getTypeStoreSize(SizeType);
Type = IRAccess::READ;
} else {
StoreInst *Store = cast<StoreInst>(Inst);
- Size = TD->getTypeStoreSize(Store->getValueOperand()->getType());
+ SizeType = Store->getValueOperand()->getType();
+ Size = TD->getTypeStoreSize(SizeType);
Type = IRAccess::WRITE;
}
@@ -159,11 +167,36 @@ IRAccess TempScopInfo::buildIRAccess(Instruction *Inst, Loop *L, Region *R) {
assert(BasePointer && "Could not find base pointer");
AccessFunction = SE->getMinusSCEV(AccessFunction, BasePointer);
+ SmallVector<const SCEV *, 4> Subscripts, Sizes;
+
+ bool IsAffine = true;
+ const SCEVAddRecExpr *AF = dyn_cast<SCEVAddRecExpr>(AccessFunction);
+
+ if (PollyDelinearize && AF) {
+ const SCEV *Remainder = AF->delinearize(*SE, Subscripts, Sizes);
+ int NSubs = Subscripts.size();
+
+ // Normalize the last dimension: integrate the size of the "scalar dimension"
+ // and the remainder of the delinearization.
+ Subscripts[NSubs-1] = SE->getMulExpr(Subscripts[NSubs-1],
+ Sizes[NSubs-1]);
+ Subscripts[NSubs-1] = SE->getAddExpr(Subscripts[NSubs-1], Remainder);
+
+ for (int i = 0; i < NSubs; ++i)
+ if (!isAffineExpr(R, Subscripts[i], *SE, BasePointer->getValue())) {
+ IsAffine = false;
+ break;
+ }
+ }
- bool IsAffine = isAffineExpr(R, AccessFunction, *SE, BasePointer->getValue());
+ if (Subscripts.size() == 0) {
+ Subscripts.push_back(AccessFunction);
+ Sizes.push_back(SE->getConstant(ZeroOffset->getType(), Size));
+ IsAffine = isAffineExpr(R, AccessFunction, *SE, BasePointer->getValue());
+ }
return IRAccess(Type, BasePointer->getValue(), AccessFunction, Size,
- IsAffine);
+ IsAffine, Subscripts, Sizes);
}
void TempScopInfo::buildAccessFunctions(Region &R, BasicBlock &BB) {
@@ -178,7 +211,11 @@ void TempScopInfo::buildAccessFunctions(Region &R, BasicBlock &BB) {
if (!isa<StoreInst>(Inst) && buildScalarDependences(Inst, &R)) {
// If the Instruction is used outside the statement, we need to build the
// write access.
- IRAccess ScalarAccess(IRAccess::SCALARWRITE, Inst, ZeroOffset, 1, true);
+ SmallVector<const SCEV *, 4> Subscripts, Sizes;
+ Subscripts.push_back(SE->getConstant(ZeroOffset->getType(), 0));
+ Sizes.push_back(SE->getConstant(ZeroOffset->getType(), 1));
+ IRAccess ScalarAccess(IRAccess::SCALARWRITE, Inst, ZeroOffset, 1, true,
+ Subscripts, Sizes);
Functions.push_back(std::make_pair(ScalarAccess, Inst));
}
}
diff --git a/polly/test/Dependences/do_pluto_matmult.ll b/polly/test/Dependences/do_pluto_matmult.ll
index 0a5d8a5cfd2..7d9799e3ed8 100644
--- a/polly/test/Dependences/do_pluto_matmult.ll
+++ b/polly/test/Dependences/do_pluto_matmult.ll
@@ -1,5 +1,5 @@
; RUN: opt %loadPolly -basicaa -polly-dependences -analyze -polly-dependences-analysis-type=value-based < %s | FileCheck %s -check-prefix=VALUE
-; RUN: opt %loadPolly -basicaa -polly-dependences -analyze -polly-dependences-analysis-type=memory-based < %s | FileCheck %s -check-prefix=MEMORY
+; RUN: opt %loadPolly -basicaa -polly-dependences -analyze -polly-dependences-analysis-type=memory-based -polly-delinearize < %s | FileCheck %s -check-prefix=MEMORY
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64"
target triple = "x86_64-unknown-linux-gnu"
diff --git a/polly/test/Isl/Ast/OpenMP/nested_loop_both_parallel_parametric.ll b/polly/test/Isl/Ast/OpenMP/nested_loop_both_parallel_parametric.ll
index 1864937f382..2c11fc5d203 100644
--- a/polly/test/Isl/Ast/OpenMP/nested_loop_both_parallel_parametric.ll
+++ b/polly/test/Isl/Ast/OpenMP/nested_loop_both_parallel_parametric.ll
@@ -1,4 +1,4 @@
-; RUN: opt %loadPolly -polly-ast -polly-ast-detect-parallel -analyze < %s | FileCheck %s
+; RUN: opt %loadPolly -polly-ast -polly-ast-detect-parallel -analyze -polly-delinearize < %s | FileCheck %s
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64"
target triple = "x86_64-pc-linux-gnu"
@@ -49,8 +49,8 @@ ret:
; to polly. However, we should be able to obtain it due to the out of bounds
; memory accesses, that would happen if n >= 1024.
;
+; CHECK: #pragma omp parallel for
; CHECK: for (int c1 = 0; c1 < n; c1 += 1)
; CHECK: #pragma simd
-; CHECK: #pragma omp parallel for
; CHECK: for (int c3 = 0; c3 < n; c3 += 1)
; CHECK: Stmt_loop_body(c1, c3);
diff --git a/polly/test/ScopInfo/constant_start_integer.ll b/polly/test/ScopInfo/constant_start_integer.ll
index 838624d69bb..5e1c901dc2e 100644
--- a/polly/test/ScopInfo/constant_start_integer.ll
+++ b/polly/test/ScopInfo/constant_start_integer.ll
@@ -1,4 +1,4 @@
-; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
+; RUN: opt %loadPolly -polly-scops -analyze -polly-delinearize < %s | FileCheck %s
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"
@@ -11,20 +11,15 @@ target triple = "x86_64-unknown-linux-gnu"
; }
; }
; }
-;
-; Access functions:
-;
-; input[j * 64 + i + 1] => {4,+,256}<%for.cond1.preheader>
-; input[j * 64 + i + 0] => {0,+,256}<%for.cond1.preheader>
-;
-; They should share the same zero-start parameter:
-;
-; p0: {0,+,256}<%for.cond1.preheader>
-; input[j * 64 + i + 1] => p0 + 4
-; input[j * 64 + i + 0] => p0
-;
-; Function Attrs: nounwind
+; CHECK p0: {0,+,256}<%for.cond1.preheader>
+; CHECK-NOT: p1
+
+; CHECK: ReadAccess
+; CHECK: [p_0] -> { Stmt_for_body3[i0] -> MemRef_input[p_0, 1 + i0] };
+; CHECK: MustWriteAccess
+; CHECK: [p_0] -> { Stmt_for_body3[i0] -> MemRef_input[p_0, i0] };
+
define void @foo(float* nocapture %input) {
entry:
br label %for.cond1.preheader
@@ -56,11 +51,3 @@ for.inc10: ; preds = %for.body3
for.end12: ; preds = %for.inc10
ret void
}
-
-; CHECK p0: {0,+,256}<%for.cond1.preheader>
-; CHECK-NOT: p1
-
-; CHECK: ReadAccess :=
-; CHECK: [p_0] -> { Stmt_for_body3[i0] -> MemRef_input[o0] : 4o0 = 4 + p_0 + 4i0 };
-; CHECK: MustWriteAccess :=
-; CHECK: [p_0] -> { Stmt_for_body3[i0] -> MemRef_input[o0] : 4o0 = p_0 + 4i0 };
diff --git a/polly/test/ScopInfo/multidim_ivs_and_integer_offsets_3d.ll b/polly/test/ScopInfo/multidim_ivs_and_integer_offsets_3d.ll
index feaae8cf002..483bdec763f 100644
--- a/polly/test/ScopInfo/multidim_ivs_and_integer_offsets_3d.ll
+++ b/polly/test/ScopInfo/multidim_ivs_and_integer_offsets_3d.ll
@@ -1,4 +1,4 @@
-; RUN: opt %loadPolly -polly-scops -analyze -polly-allow-nonaffine < %s | FileCheck %s
+; RUN: opt %loadPolly -polly-scops -analyze -polly-delinearize < %s | FileCheck %s
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"
@@ -9,14 +9,18 @@ target triple = "x86_64-unknown-linux-gnu"
; for (long k = 0; k < o; k++)
; A[i+3][j-4][k+7] = 1.0;
; }
-;
-; Access function:
-;
-; {{{(56 + (8 * (-4 + (3 * %m)) * %o) + %A),+,(8 * %m * %o)}<%for.i>,+,
-; (8 * %o)}<%for.j>,+,8}<%for.k>
-;
-; TODO: Recover the multi-dimensional array information to avoid the
-; conservative approximation we are using today.
+
+; CHECK: p0: %n
+; CHECK: p1: %m
+; CHECK: p2: %o
+; CHECK-NOT: p3
+
+; CHECK: Domain
+; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] : i0 >= 0 and i0 <= -1 + n and i1 >= 0 and i1 <= -1 + m and i2 >= 0 and i2 <= -1 + o };
+; CHECK: Scattering
+; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> scattering[0, i0, 0, i1, 0, i2, 0] };
+; CHECK: MustWriteAccess
+; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> MemRef_A[3 + i0, -4 + i1, 7 + i2] };
define void @foo(i64 %n, i64 %m, i64 %o, double* %A) {
entry:
@@ -61,16 +65,3 @@ for.i.inc:
end:
ret void
}
-
-; CHECK: p0: %n
-; CHECK: p1: %m
-; CHECK: p2: %o
-; CHECK-NOT: p3
-
-; CHECK: Domain
-; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] : i0 >= 0 and i0 <= -1 + n and i1 >= 0 and i1 <= -1 + m and i2 >= 0 and i2 <= -1 + o };
-; CHECK: Scattering
-; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> scattering[0, i0, 0, i1, 0, i2, 0] };
-; CHECK: MayWriteAccess
-; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> MemRef_A[o0] };
-
diff --git a/polly/test/ScopInfo/multidim_ivs_and_parameteric_offsets_3d.ll b/polly/test/ScopInfo/multidim_ivs_and_parameteric_offsets_3d.ll
index 635ba9b7fab..c32a8dbf663 100644
--- a/polly/test/ScopInfo/multidim_ivs_and_parameteric_offsets_3d.ll
+++ b/polly/test/ScopInfo/multidim_ivs_and_parameteric_offsets_3d.ll
@@ -1,4 +1,4 @@
-; RUN: opt %loadPolly -polly-scops -analyze -polly-allow-nonaffine < %s | FileCheck %s
+; RUN: opt %loadPolly -polly-scops -analyze -polly-delinearize < %s | FileCheck %s
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"
@@ -14,8 +14,20 @@ target triple = "x86_64-unknown-linux-gnu"
; {{{((8 * ((((%m * %p) + %q) * %o) + %r)) + %A),+,(8 * %m * %o)}<%for.i>,+,
; (8 * %o)}<%for.j>,+,8}<%for.k>
;
-; TODO: Recover the multi-dimensional array information to avoid the
-; conservative approximation we are using today.
+; CHECK: p0: %n
+; CHECK: p1: %m
+; CHECK: p2: %o
+; CHECK: p3: %p
+; CHECK: p4: %q
+; CHECK: p5: %r
+; CHECK-NOT: p6
+;
+; CHECK: Domain
+; CHECK: [n, m, o, p, q, r] -> { Stmt_for_k[i0, i1, i2] : i0 >= 0 and i0 <= -1 + n and i1 >= 0 and i1 <= -1 + m and i2 >= 0 and i2 <= -1 + o };
+; CHECK: Scattering
+; CHECK: [n, m, o, p, q, r] -> { Stmt_for_k[i0, i1, i2] -> scattering[0, i0, 0, i1, 0, i2, 0] };
+; CHECK: MustWriteAccess
+; CHECK: [n, m, o, p, q, r] -> { Stmt_for_k[i0, i1, i2] -> MemRef_A[p + i0, q + i1, r + i2] };
define void @foo(i64 %n, i64 %m, i64 %o, double* %A, i64 %p, i64 %q, i64 %r) {
entry:
@@ -60,15 +72,3 @@ for.i.inc:
end:
ret void
}
-
-; CHECK: p0: %n
-; CHECK: p1: %m
-; CHECK: p2: %o
-; CHECK-NOT: p3
-
-; CHECK: Domain
-; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] : i0 >= 0 and i0 <= -1 + n and i1 >= 0 and i1 <= -1 + m and i2 >= 0 and i2 <= -1 + o };
-; CHECK: Scattering
-; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> scattering[0, i0, 0, i1, 0, i2, 0] };
-; CHECK: MayWriteAccess
-; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> MemRef_A[o0] };
diff --git a/polly/test/ScopInfo/multidim_nested_start_integer.ll b/polly/test/ScopInfo/multidim_nested_start_integer.ll
index dd6bcd6e725..18c6a21a306 100644
--- a/polly/test/ScopInfo/multidim_nested_start_integer.ll
+++ b/polly/test/ScopInfo/multidim_nested_start_integer.ll
@@ -1,4 +1,4 @@
-; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
+; RUN: opt %loadPolly -polly-scops -analyze -polly-delinearize < %s | FileCheck %s
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"
@@ -14,11 +14,17 @@ target triple = "x86_64-unknown-linux-gnu"
; {{{(56 + (8 * (-4 + (3 * %m)) * %o) + %A),+,(8 * %m * %o)}<%for.i>,+,
; (8 * %o)}<%for.j>,+,8}<%for.k>
;
-; The nested 'start' should be splitted into three parameters:
-; p1: {0,+,(8 * %o)}<%for.j>
-; p2: {0,+,(8 * %m * %o)}<%for.i>
-; p3: (8 * (-4 + (3 * %m)) * %o)
-;
+; CHECK: p0: %n
+; CHECK: p1: %m
+; CHECK: p2: %o
+; CHECK-NOT: p3
+; CHECK: Domain
+; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] : i0 >= 0 and i0 <= -1 + n and i1 >= 0 and i1 <= -1 + m and i2 >= 0 and i2 <= -1 + o };
+; CHECK: Scattering
+; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> scattering[0, i0, 0, i1, 0, i2, 0] };
+; CHECK: MustWriteAccess
+; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> MemRef_A[3 + i0, -4 + i1, 7 + i2] };
+
define void @foo(i64 %n, i64 %m, i64 %o, double* %A) {
entry:
@@ -63,16 +69,3 @@ for.i.inc:
end:
ret void
}
-
-; CHECK: p0: %o
-; CHECK: p1: {0,+,(8 * %o)}<%for.j>
-; CHECK: p2: {0,+,(8 * %m * %o)}<%for.i>
-; CHECK: p3: (8 * (-4 + (3 * %m)) * %o)
-; CHECK-NOT: p4
-
-; CHECK: Domain
-; CHECK: [o, p_1, p_2, p_3] -> { Stmt_for_k[i0] : i0 >= 0 and i0 <= -1 + o };
-; CHECK: Scattering
-; CHECK: [o, p_1, p_2, p_3] -> { Stmt_for_k[i0] -> scattering[0, i0, 0] };
-; CHECK: MustWriteAccess
-; CHECK: [o, p_1, p_2, p_3] -> { Stmt_for_k[i0] -> MemRef_A[o0] : 8o0 = 56 + p_1 + p_2 + p_3 + 8i0 };
diff --git a/polly/test/ScopInfo/multidim_nested_start_share_parameter.ll b/polly/test/ScopInfo/multidim_nested_start_share_parameter.ll
index d197ea0f2ba..6bcc73aa488 100644
--- a/polly/test/ScopInfo/multidim_nested_start_share_parameter.ll
+++ b/polly/test/ScopInfo/multidim_nested_start_share_parameter.ll
@@ -1,4 +1,4 @@
-; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
+; RUN: opt %loadPolly -polly-scops -analyze -polly-delinearize < %s | FileCheck %s
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"
@@ -19,10 +19,14 @@ target triple = "x86_64-unknown-linux-gnu"
; {{{(136 + (8 * (-14 + (13 * %m)) * %o) + %A),+,(8 * %m * %o)}<%for.i>,+,
; (8 * %o)}<%for.j>,+,8}<%for.k>
;
-; They should share the following parameters:
-; p1: {0,+,(8 * %o)}<%for.j>
-; p2: {0,+,(8 * %m * %o)}<%for.i>
+; CHECK: p0: %n
+; CHECK: p1: %m
+; CHECK: p2: %o
+; CHECK-NOT: p3
;
+; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> MemRef_A[3 + i0, -4 + i1, 7 + i2] };
+; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> MemRef_A[13 + i0, -14 + i1, 17 + i2] };
+
define void @foo(i64 %n, i64 %m, i64 %o, double* %A) {
entry:
@@ -79,14 +83,3 @@ for.i.inc:
end:
ret void
}
-
-; CHECK: p0: %o
-; CHECK: p1: {0,+,(8 * %o)}<%for.j>
-; CHECK: p2: {0,+,(8 * %m * %o)}<%for.i>
-; CHECK: p3: (8 * (-4 + (3 * %m)) * %o)
-; CHECK: p4: (8 * (-14 + (13 * %m)) * %o)
-; CHECK-NOT: p4
-
-; CHECK: [o, p_1, p_2, p_3, p_4] -> { Stmt_for_k[i0] -> MemRef_A[o0] : 8o0 = 56 + p_1 + p_2 + p_3 + 8i0 };
-; CHECK: [o, p_1, p_2, p_3, p_4] -> { Stmt_for_k[i0] -> MemRef_A[o0] : 8o0 = 136 + p_1 + p_2 + p_4 + 8i0 };
-
diff --git a/polly/test/ScopInfo/multidim_only_ivs_2d.ll b/polly/test/ScopInfo/multidim_only_ivs_2d.ll
index e99674292b2..b268ddbab5a 100644
--- a/polly/test/ScopInfo/multidim_only_ivs_2d.ll
+++ b/polly/test/ScopInfo/multidim_only_ivs_2d.ll
@@ -1,4 +1,4 @@
-; RUN: opt %loadPolly -polly-scops -analyze -polly-allow-nonaffine < %s | FileCheck %s
+; RUN: opt %loadPolly -polly-scops -analyze -polly-delinearize < %s | FileCheck %s
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"
@@ -9,11 +9,17 @@ target triple = "x86_64-unknown-linux-gnu"
; for (long j = 0; j < m; j++)
; A[i][j] = 1.0;
; }
-;
-; Access function: {{0,+,%m}<%for.i>,+,1}<nw><%for.j>
-;
-; TODO: Recover the multi-dimensional array information to avoid the
-; conservative approximation we are using today.
+
+; CHECK: p0: %n
+; CHECK: p1: %m
+; CHECK-NOT: p3
+
+; CHECK: Domain
+; CHECK: [n, m] -> { Stmt_for_j[i0, i1] : i0 >= 0 and i0 <= -1 + n and i1 >= 0 and i1 <= -1 + m };
+; CHECK: Scattering
+; CHECK: [n, m] -> { Stmt_for_j[i0, i1] -> scattering[0, i0, 0, i1, 0] };
+; CHECK: MustWriteAccess
+; CHECK: [n, m] -> { Stmt_for_j[i0, i1] -> MemRef_A[i0, i1] };
define void @foo(i64 %n, i64 %m, double* %A) {
entry:
@@ -41,15 +47,3 @@ for.i.inc:
end:
ret void
}
-
-; CHECK: p0: %n
-; CHECK: p1: %m
-; CHECK-NOT: p3
-
-; CHECK: Domain :=
-; CHECK: [n, m] -> { Stmt_for_j[i0, i1] : i0 >= 0 and i0 <= -1 + n and i1 >= 0 and i1 <= -1 + m };
-; CHECK: Scattering :=
-; CHECK: [n, m] -> { Stmt_for_j[i0, i1] -> scattering[0, i0, 0, i1, 0] };
-; CHECK: MayWriteAccess :=
-; CHECK: [n, m] -> { Stmt_for_j[i0, i1] -> MemRef_A[o0] };
-
diff --git a/polly/test/ScopInfo/multidim_only_ivs_3d.ll b/polly/test/ScopInfo/multidim_only_ivs_3d.ll
index 36f2a17f4e1..01b2fbd38cd 100644
--- a/polly/test/ScopInfo/multidim_only_ivs_3d.ll
+++ b/polly/test/ScopInfo/multidim_only_ivs_3d.ll
@@ -1,4 +1,4 @@
-; RUN: opt %loadPolly -polly-scops -analyze -polly-allow-nonaffine < %s | FileCheck %s
+; RUN: opt %loadPolly -polly-scops -analyze -polly-delinearize < %s | FileCheck %s
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"
@@ -9,11 +9,18 @@ target triple = "x86_64-unknown-linux-gnu"
; for (long k = 0; k < o; k++)
; A[i][j][k] = 1.0;
; }
+
+; CHECK: p0: %n
+; CHECK: p1: %m
+; CHECK: p2: %o
+; CHECK-NOT: p3
;
-; Access function: {{{0,+,(%m * %o)}<%for.i>,+,%o}<%for.j>,+,1}<nw><%for.k>
-;
-; TODO: Recover the multi-dimensional array information to avoid the
-; conservative approximation we are using today.
+; CHECK: Domain
+; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] : i0 >= 0 and i0 <= -1 + n and i1 >= 0 and i1 <= -1 + m and i2 >= 0 and i2 <= -1 + o };
+; CHECK: Scattering
+; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> scattering[0, i0, 0, i1, 0, i2, 0] };
+; CHECK: WriteAccess
+; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> MemRef_A[i0, i1, i2] };
define void @foo(i64 %n, i64 %m, i64 %o, double* %A) {
entry:
@@ -55,16 +62,3 @@ for.i.inc:
end:
ret void
}
-
-; CHECK: p0: %n
-; CHECK: p1: %m
-; CHECK: p2: %o
-; CHECK-NOT: p3
-
-; CHECK: Domain
-; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] : i0 >= 0 and i0 <= -1 + n and i1 >= 0 and i1 <= -1 + m and i2 >= 0 and i2 <= -1 + o };
-; CHECK: Scattering
-; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> scattering[0, i0, 0, i1, 0, i2, 0] };
-; CHECK: WriteAccess
-; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> MemRef_A[o0] };
-
diff --git a/polly/test/ScopInfo/multidim_only_ivs_3d_cast.ll b/polly/test/ScopInfo/multidim_only_ivs_3d_cast.ll
index bcc1a5062dc..b4c27ce50b0 100644
--- a/polly/test/ScopInfo/multidim_only_ivs_3d_cast.ll
+++ b/polly/test/ScopInfo/multidim_only_ivs_3d_cast.ll
@@ -1,4 +1,4 @@
-; RUN: opt %loadPolly -polly-scops -analyze -polly-allow-nonaffine < %s | FileCheck %s
+; RUN: opt %loadPolly -polly-scops -analyze -polly-delinearize < %s | FileCheck %s
; void foo(int n, int m, int o, double A[n][m][o]) {
;
@@ -7,13 +7,18 @@
; for (int k = 0; k < o; k++)
; A[i][j][k] = 1.0;
; }
-;
-; Access function:
-; {{{%A,+,(8 * (zext i32 %m to i64) * (zext i32 %o to i64))}<%for.i>,+,
-; (8 * (zext i32 %o to i64))}<%for.j>,+,8}<%for.k>
-;
-; TODO: Recover the multi-dimensional array information to avoid the
-; conservative approximation we are using today.
+
+; CHECK: p0: %n
+; CHECK: p1: %m
+; CHECK: p2: %o
+; CHECK-NOT: p3
+
+; CHECK: Domain
+; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] : i0 >= 0 and i0 <= -1 + n and i1 >= 0 and i1 <= -1 + m and i2 >= 0 and i2 <= -1 + o };
+; CHECK: Scattering
+; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> scattering[0, i0, 0, i1, 0, i2, 0] };
+; CHECK: WriteAccess
+; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> MemRef_A[i0, i1, i2] };
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"
@@ -66,15 +71,3 @@ for.i.inc:
end:
ret void
}
-
-; CHECK: p0: %n
-; CHECK: p1: %m
-; CHECK: p2: %o
-; CHECK-NOT: p3
-
-; CHECK: Domain
-; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] : i0 >= 0 and i0 <= -1 + n and i1 >= 0 and i1 <= -1 + m and i2 >= 0 and i2 <= -1 + o };
-; CHECK: Scattering
-; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> scattering[0, i0, 0, i1, 0, i2, 0] };
-; CHECK: WriteAccess
-; CHECK: [n, m, o] -> { Stmt_for_k[i0, i1, i2] -> MemRef_A[o0]
OpenPOWER on IntegriCloud