Index: lib/Transforms/Scalar/SeparateConstOffsetFromGEP.cpp =================================================================== --- lib/Transforms/Scalar/SeparateConstOffsetFromGEP.cpp +++ lib/Transforms/Scalar/SeparateConstOffsetFromGEP.cpp @@ -236,6 +236,15 @@ AU.addRequired(); AU.addRequired(); } + + bool doInitialization(Module &M) override { + DataLayoutPass *DLP = getAnalysisIfAvailable(); + if (DLP == nullptr) + report_fatal_error("data layout missing"); + DL = &DLP->getDataLayout(); + return false; + } + bool runOnFunction(Function &F) override; private: @@ -246,8 +255,42 @@ /// function only inspects the GEP without changing it. The output /// NeedsExtraction indicates whether we can extract a non-zero constant /// offset from any index. - int64_t accumulateByteOffset(GetElementPtrInst *GEP, const DataLayout *DL, - bool &NeedsExtraction); + int64_t accumulateByteOffset(GetElementPtrInst *GEP, bool &NeedsExtraction); + /// Canonicalize array indices to pointer-size integers. This helps to + /// simplify the logic of splitting a GEP. For example, if a + b is a + /// pointer-size integer, we have + /// gep base, a + b = gep (gep base, a), b + /// However, this equality may not hold if the size of a + b is smaller than + /// the pointer size, because LLVM conceptually sign-extends GEP indices to + /// pointer size before computing the address + /// (http://llvm.org/docs/LangRef.html#id181). + /// + /// 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; }; } // anonymous namespace @@ -297,9 +340,9 @@ // 1 | 0 | sext(BO) == sext(A) op sext(B) // 1 | 1 | zext(sext(BO)) == // | | zext(sext(A)) op zext(sext(B)) - if (BO->getOpcode() == Instruction::Add && NonNegative) { + if (BO->getOpcode() == Instruction::Add && !ZeroExtended && NonNegative) { // If a + b >= 0 and (a >= 0 or b >= 0), then - // s/zext(a + b) = s/zext(a) + s/zext(b) + // sext(a + b) = sext(a) + sext(b) // even if the addition is not marked nsw. // // Leveraging this invarient, we can trace into an sext'ed inbound GEP @@ -379,7 +422,6 @@ // 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); @@ -553,8 +595,64 @@ return (LHSKnownZero | RHSKnownZero).isAllOnesValue(); } -int64_t SeparateConstOffsetFromGEP::accumulateByteOffset( - GetElementPtrInst *GEP, const DataLayout *DL, bool &NeedsExtraction) { +bool SeparateConstOffsetFromGEP::canonicalizeArrayIndicesToPointerSize( + GetElementPtrInst *GEP) { + bool Changed = false; + Type *IntPtrTy = DL->getIntPtrType(GEP->getType()); + gep_type_iterator GTI = gep_type_begin(*GEP); + for (User::op_iterator I = GEP->op_begin() + 1, E = GEP->op_end(); + I != E; ++I, ++GTI) { + // Skip struct member indices which must be i32. + if (isa(*GTI)) { + if ((*I)->getType() != IntPtrTy) { + *I = CastInst::CreateIntegerCast(*I, IntPtrTy, true, "idxprom", GEP); + Changed = true; + } + } + } + return Changed; +} + +bool +SeparateConstOffsetFromGEP::convertInBoundsZExtToSExt(GetElementPtrInst *GEP) { + if (!GEP->isInBounds()) + return false; + + // TODO: consider alloca + GlobalVariable *UnderlyingObject = + dyn_cast(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(*GTI)) { + if (ZExtInst *Extended = dyn_cast(*I)) { + unsigned SrcBitWidth = + cast(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) { NeedsExtraction = false; int64_t AccumulativeByteOffset = 0; gep_type_iterator GTI = gep_type_begin(*GEP); @@ -587,35 +685,11 @@ return false; bool Changed = false; - // Canonicalize array indices to pointer-size integers. This helps to simplify - // the logic of splitting a GEP. For example, if a + b is a pointer-size - // integer, we have - // gep base, a + b = gep (gep base, a), b - // However, this equality may not hold if the size of a + b is smaller than - // the pointer size, because LLVM conceptually sign-extends GEP indices to - // pointer size before computing the address - // (http://llvm.org/docs/LangRef.html#id181). - // - // This canonicalization is very likely already done in clang and instcombine. - // Therefore, the program will probably remain the same. - // - // Verified in @i32_add in split-gep.ll - const DataLayout *DL = &getAnalysis().getDataLayout(); - Type *IntPtrTy = DL->getIntPtrType(GEP->getType()); - gep_type_iterator GTI = gep_type_begin(*GEP); - for (User::op_iterator I = GEP->op_begin() + 1, E = GEP->op_end(); - I != E; ++I, ++GTI) { - if (isa(*GTI)) { - if ((*I)->getType() != IntPtrTy) { - *I = CastInst::CreateIntegerCast(*I, IntPtrTy, true, "idxprom", GEP); - Changed = true; - } - } - } + Changed |= canonicalizeArrayIndicesToPointerSize(GEP); + Changed |= convertInBoundsZExtToSExt(GEP); bool NeedsExtraction; - int64_t AccumulativeByteOffset = - accumulateByteOffset(GEP, DL, NeedsExtraction); + int64_t AccumulativeByteOffset = accumulateByteOffset(GEP, NeedsExtraction); if (!NeedsExtraction) return Changed; @@ -631,7 +705,7 @@ // Remove the constant offset in each GEP index. The resultant GEP computes // the variadic base. - GTI = gep_type_begin(*GEP); + gep_type_iterator GTI = gep_type_begin(*GEP); for (unsigned I = 1, E = GEP->getNumOperands(); I != E; ++I, ++GTI) { if (isa(*GTI)) { Value *NewIdx = nullptr; @@ -699,6 +773,7 @@ uint64_t ElementTypeSizeOfGEP = DL->getTypeAllocSize(GEP->getType()->getElementType()); + Type *IntPtrTy = DL->getIntPtrType(GEP->getType()); if (AccumulativeByteOffset % ElementTypeSizeOfGEP == 0) { // Very likely. As long as %gep is natually aligned, the byte offset we // extracted should be a multiple of sizeof(*%gep). Index: test/Transforms/SeparateConstOffsetFromGEP/NVPTX/split-gep-and-gvn.ll =================================================================== --- test/Transforms/SeparateConstOffsetFromGEP/NVPTX/split-gep-and-gvn.ll +++ test/Transforms/SeparateConstOffsetFromGEP/NVPTX/split-gep-and-gvn.ll @@ -98,3 +98,44 @@ ; 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 Index: test/Transforms/SeparateConstOffsetFromGEP/NVPTX/split-gep.ll =================================================================== --- test/Transforms/SeparateConstOffsetFromGEP/NVPTX/split-gep.ll +++ test/Transforms/SeparateConstOffsetFromGEP/NVPTX/split-gep.ll @@ -26,21 +26,23 @@ ; 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 @@ -73,16 +75,16 @@ %2 = zext i48 %1 to i64 ; zext(sext(a +nsw nuw 1)) = zext(sext(a)) + 1 %3 = add nsw i32 %b, 2 %4 = sext i32 %3 to i48 - %5 = zext i48 %4 to i64 ; zext(sext(a +nsw 2)) != zext(sext(a)) + 2 - %p1 = getelementptr inbounds [32 x [32 x float]]* @float_2d_array, i64 0, i64 %2, i64 %5 + %5 = zext i48 %4 to i64 ; zext(sext(b +nsw 2)) != zext(sext(b)) + 2 + %p1 = getelementptr [32 x [32 x float]]* @float_2d_array, i64 0, i64 %2, i64 %5 store float* %p1, float** %out1 %6 = add nuw i32 %a, 3 %7 = zext i32 %6 to i48 - %8 = sext i48 %7 to i64 ; sext(zext(b +nuw 3)) = zext(b +nuw 3) = zext(b) + 3 + %8 = sext i48 %7 to i64 ; sext(zext(a +nuw 3)) = zext(a +nuw 3) = zext(a) + 3 %9 = add nsw i32 %b, 4 %10 = zext i32 %9 to i48 %11 = sext i48 %10 to i64 ; sext(zext(b +nsw 4)) != zext(b) + 4 - %p2 = getelementptr inbounds [32 x [32 x float]]* @float_2d_array, i64 0, i64 %8, i64 %11 + %p2 = getelementptr [32 x [32 x float]]* @float_2d_array, i64 0, i64 %8, i64 %11 store float* %p2, float** %out2 ret void } @@ -232,3 +234,28 @@ ; 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