maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

delinearize memory access functions 

llvm-svn: 205799
This commit is contained in:
Sebastian Pop 2014-04-08 21:20:44 +00:00
parent 0a31e04219
commit 1801668af3
15 changed files with 201 additions and 175 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
polly/include/polly/ScopDetection.h
View File

@ -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

10
polly/include/polly/TempScopInfo.h
View File

@ -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; }

25
polly/lib/Analysis/ScopDetection.cpp
View File

@ -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);
const SCEVAddRecExpr *AF = dyn_cast<SCEVAddRecExpr>(AccessFunction);
if (!AllowNonAffine &&
!isAffineExpr(&Context.CurRegion, AccessFunction, *SE, BaseValue))
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.

42
polly/lib/Analysis/ScopInfo.cpp
View File

@ -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;
// 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);
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);
}
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));

49
polly/lib/Analysis/TempScopInfo.cpp
View File

@ -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 = isAffineExpr(R, AccessFunction, *SE, BasePointer->getValue());
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;
}
}
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));
}
}

2
polly/test/Dependences/do_pluto_matmult.ll
View File

@ -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"

4
polly/test/Isl/Ast/OpenMP/nested_loop_both_parallel_parametric.ll
View File

@ -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);

31
polly/test/ScopInfo/constant_start_integer.ll
View File

@ -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 };

35
polly/test/ScopInfo/multidim_ivs_and_integer_offsets_3d.ll
View File

@ -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] };

30
polly/test/ScopInfo/multidim_ivs_and_parameteric_offsets_3d.ll
View File

@ -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] };

31
polly/test/ScopInfo/multidim_nested_start_integer.ll
View File

@ -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 };

23
polly/test/ScopInfo/multidim_nested_start_share_parameter.ll
View File

@ -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 };

30
polly/test/ScopInfo/multidim_only_ivs_2d.ll
View File

@ -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] };

30
polly/test/ScopInfo/multidim_only_ivs_3d.ll
View File

@ -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] };

33
polly/test/ScopInfo/multidim_only_ivs_3d_cast.ll
View File

@ -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]