6502/llvm-6502
1
0
Fork 0
mirror of https://github.com/c64scene-ar/llvm-6502.git synced 2024-12-12 13:30:51 +00:00
Code Issues Projects Releases Wiki Activity

[SeparateConstOffsetFromGEP] inbounds zext => sext for better splitting

Browse Source 
For each array index that is in the form of zext(a), convert it to sext(a)
if we can prove zext(a) <= max signed value of typeof(a). The conversion
helps to split zext(x + y) into sext(x) + sext(y).

Reviewed in http://reviews.llvm.org/D4060



git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@210444 91177308-0d34-0410-b5e6-96231b3b80d8
This commit is contained in:
Jingyue Wu 2014-06-08 23:49:34 +00:00
parent 66a3b9049b
commit 09d6a9040c
3 changed files with 128 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

58
lib/Transforms/Scalar/SeparateConstOffsetFromGEP.cpp
View File

@ -268,8 +268,27 @@ class SeparateConstOffsetFromGEP : public FunctionPass {
/// This canonicalization is very likely already done in clang and
/// instcombine. Therefore, the program will probably remain the same.
///
/// Returns true if the module changes.
///
/// Verified in @i32_add in split-gep.ll
bool canonicalizeArrayIndicesToPointerSize(GetElementPtrInst *GEP);
/// For each array index that is in the form of zext(a), convert it to sext(a)
/// if we can prove zext(a) <= max signed value of typeof(a). We prefer
/// sext(a) to zext(a), because in the special case where x + y >= 0 and
/// (x >= 0 or y >= 0), function CanTraceInto can split sext(x + y),
/// while no such case exists for zext(x + y).
///
/// Note that
/// zext(x + y) = zext(x) + zext(y)
/// is wrong, e.g.,
/// zext i32(UINT_MAX + 1) to i64 !=
/// (zext i32 UINT_MAX to i64) + (zext i32 1 to i64)
///
/// Returns true if the module changes.
///
/// Verified in @inbounds_zext_add in split-gep.ll and @sum_of_array3 in
/// split-gep-and-gvn.ll
bool convertInBoundsZExtToSExt(GetElementPtrInst *GEP);
const DataLayout *DL;
};
@ -403,7 +422,6 @@ APInt ConstantOffsetExtractor::find(Value *V, bool SignExtended,
// sext(zext(a)) = zext(a). Verified in @sext_zext in split-gep.ll.
//
// Clear the NonNegative flag, because zext(a) >= 0 does not imply a >= 0.
// TODO: if zext(a) < 2 ^ (bitwidth(a) - 1), we can prove a >= 0.
ConstantOffset =
find(U->getOperand(0), /* SignExtended */ false,
/* ZeroExtended */ true, /* NonNegative */ false).zext(BitWidth);
@ -595,6 +613,43 @@ bool SeparateConstOffsetFromGEP::canonicalizeArrayIndicesToPointerSize(
return Changed;
}
bool
SeparateConstOffsetFromGEP::convertInBoundsZExtToSExt(GetElementPtrInst *GEP) {
if (!GEP->isInBounds())
return false;
// TODO: consider alloca
GlobalVariable *UnderlyingObject =
dyn_cast<GlobalVariable>(GEP->getPointerOperand());
if (UnderlyingObject == nullptr)
return false;
uint64_t ObjectSize =
DL->getTypeAllocSize(UnderlyingObject->getType()->getElementType());
gep_type_iterator GTI = gep_type_begin(*GEP);
bool Changed = false;
for (User::op_iterator I = GEP->op_begin() + 1, E = GEP->op_end(); I != E;
++I, ++GTI) {
if (isa<SequentialType>(*GTI)) {
if (ZExtInst *Extended = dyn_cast<ZExtInst>(*I)) {
unsigned SrcBitWidth =
cast<IntegerType>(Extended->getSrcTy())->getBitWidth();
// For GEP operand zext(a), if a <= max signed value of typeof(a), then
// the sign bit of a is zero and sext(a) = zext(a). Because the GEP is
// in bounds, we know a <= ObjectSize, so the condition can be reduced
// to ObjectSize <= max signed value of typeof(a).
if (ObjectSize <=
APInt::getSignedMaxValue(SrcBitWidth).getZExtValue()) {
*I = new SExtInst(Extended->getOperand(0), Extended->getType(),
Extended->getName(), GEP);
Changed = true;
}
}
}
}
return Changed;
}
int64_t
SeparateConstOffsetFromGEP::accumulateByteOffset(GetElementPtrInst *GEP,
bool &NeedsExtraction) {
@ -631,6 +686,7 @@ bool SeparateConstOffsetFromGEP::splitGEP(GetElementPtrInst *GEP) {
bool Changed = false;
Changed |= canonicalizeArrayIndicesToPointerSize(GEP);
Changed |= convertInBoundsZExtToSExt(GEP);
bool NeedsExtraction;
int64_t AccumulativeByteOffset = accumulateByteOffset(GEP, NeedsExtraction);

41
test/Transforms/SeparateConstOffsetFromGEP/NVPTX/split-gep-and-gvn.ll
View File

@ -98,3 +98,44 @@ define void @sum_of_array2(i32 %x, i32 %y, float* nocapture %output) {
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 1
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 32
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 33
; Similar to @sum_of_array3, but extends array indices using zext instead of
; sext. e.g., array[zext(x + 1)][zext(y + 1)].
define void @sum_of_array3(i32 %x, i32 %y, float* nocapture %output) {
.preheader:
%0 = zext i32 %y to i64
%1 = zext i32 %x to i64
%2 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %0
%3 = addrspacecast float addrspace(3)* %2 to float*
%4 = load float* %3, align 4
%5 = fadd float %4, 0.000000e+00
%6 = add i32 %y, 1
%7 = zext i32 %6 to i64
%8 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %7
%9 = addrspacecast float addrspace(3)* %8 to float*
%10 = load float* %9, align 4
%11 = fadd float %5, %10
%12 = add i32 %x, 1
%13 = zext i32 %12 to i64
%14 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %13, i64 %0
%15 = addrspacecast float addrspace(3)* %14 to float*
%16 = load float* %15, align 4
%17 = fadd float %11, %16
%18 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %13, i64 %7
%19 = addrspacecast float addrspace(3)* %18 to float*
%20 = load float* %19, align 4
%21 = fadd float %17, %20
store float %21, float* %output, align 4
ret void
}
; PTX-LABEL: sum_of_array3(
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG:%(rl|r)[0-9]+]]{{\]}}
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+4{{\]}}
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+128{{\]}}
; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+132{{\]}}
; IR-LABEL: @sum_of_array3(
; IR: [[BASE_PTR:%[a-zA-Z0-9]+]] = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}}
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 1
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 32
; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 33

33
test/Transforms/SeparateConstOffsetFromGEP/NVPTX/split-gep.ll
View File

@ -26,21 +26,23 @@ entry:
; CHECK-LABEL: @struct(
; CHECK: getelementptr [1024 x %struct.S]* @struct_array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i32 1
; We should be able to trace into s/zext(a + b) if a + b is non-negative
; We should be able to trace into sext(a + b) if a + b is non-negative
; (e.g., used as an index of an inbounds GEP) and one of a and b is
; non-negative.
define float* @sext_add(i32 %i, i32 %j) {
entry:
%0 = add i32 %i, 1
%1 = sext i32 %0 to i64 ; inbound sext(i + 1) = sext(i) + 1
%2 = sub i32 %j, 2
; However, inbound sext(j - 2) != sext(j) - 2, e.g., j = INT_MIN
%2 = add i32 %j, -2
; However, inbound sext(j + -2) != sext(j) + -2, e.g., j = INT_MIN
%3 = sext i32 %2 to i64
%p = getelementptr inbounds [32 x [32 x float]]* @float_2d_array, i64 0, i64 %1, i64 %3
ret float* %p
}
; CHECK-LABEL: @sext_add(
; CHECK-NOT: = add
; CHECK: add i32 %j, -2
; CHECK: sext
; CHECK: getelementptr [32 x [32 x float]]* @float_2d_array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}}
; CHECK: getelementptr float* %{{[a-zA-Z0-9]+}}, i64 32
@ -232,3 +234,28 @@ entry:
; CHECK-LABEL: @and(
; CHECK: getelementptr [32 x [32 x float]]* @float_2d_array
; CHECK-NOT: getelementptr
; if zext(a + b) <= max signed value of typeof(a + b), then we can prove
; a + b >= 0 and zext(a + b) == sext(a + b). If we can prove further a or b is
; non-negative, we have zext(a + b) == sext(a) + sext(b).
define float* @inbounds_zext_add(i32 %i, i4 %j) {
entry:
%0 = add i32 %i, 1
%1 = zext i32 %0 to i64
; Because zext(i + 1) is an index of an in bounds GEP based on
; float_2d_array, zext(i + 1) <= sizeof(float_2d_array) = 4096.
; Furthermore, since typeof(i + 1) is i32 and 4096 < 2^31, we are sure the
; sign bit of i + 1 is 0. This implies zext(i + 1) = sext(i + 1).
%2 = add i4 %j, 2
%3 = zext i4 %2 to i64
; In this case, typeof(j + 2) is i4, so zext(j + 2) <= 4096 does not imply
; the sign bit of j + 2 is 0.
%p = getelementptr inbounds [32 x [32 x float]]* @float_2d_array, i64 0, i64 %1, i64 %3
ret float* %p
}
; CHECK-LABEL: @inbounds_zext_add(
; CHECK-NOT: add
; CHECK: add i4 %j, 2
; CHECK: sext
; CHECK: getelementptr [32 x [32 x float]]* @float_2d_array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}}
; CHECK: getelementptr float* %{{[a-zA-Z0-9]+}}, i64 32