//===-- SeparateConstOffsetFromGEP.cpp - ------------------------*- C++ -*-===// // // The LLVM Compiler Infrastructure // // This file is distributed under the University of Illinois Open Source // License. See LICENSE.TXT for details. // //===----------------------------------------------------------------------===// // // Loop unrolling may create many similar GEPs for array accesses. // e.g., a 2-level loop // // float a[32][32]; // global variable // // for (int i = 0; i < 2; ++i) { // for (int j = 0; j < 2; ++j) { // ... // ... = a[x + i][y + j]; // ... // } // } // // will probably be unrolled to: // // gep %a, 0, %x, %y; load // gep %a, 0, %x, %y + 1; load // gep %a, 0, %x + 1, %y; load // gep %a, 0, %x + 1, %y + 1; load // // LLVM's GVN does not use partial redundancy elimination yet, and is thus // unable to reuse (gep %a, 0, %x, %y). As a result, this misoptimization incurs // significant slowdown in targets with limited addressing modes. For instance, // because the PTX target does not support the reg+reg addressing mode, the // NVPTX backend emits PTX code that literally computes the pointer address of // each GEP, wasting tons of registers. It emits the following PTX for the // first load and similar PTX for other loads. // // mov.u32 %r1, %x; // mov.u32 %r2, %y; // mul.wide.u32 %rl2, %r1, 128; // mov.u64 %rl3, a; // add.s64 %rl4, %rl3, %rl2; // mul.wide.u32 %rl5, %r2, 4; // add.s64 %rl6, %rl4, %rl5; // ld.global.f32 %f1, [%rl6]; // // To reduce the register pressure, the optimization implemented in this file // merges the common part of a group of GEPs, so we can compute each pointer // address by adding a simple offset to the common part, saving many registers. // // It works by splitting each GEP into a variadic base and a constant offset. // The variadic base can be computed once and reused by multiple GEPs, and the // constant offsets can be nicely folded into the reg+immediate addressing mode // (supported by most targets) without using any extra register. // // For instance, we transform the four GEPs and four loads in the above example // into: // // base = gep a, 0, x, y // load base // laod base + 1 * sizeof(float) // load base + 32 * sizeof(float) // load base + 33 * sizeof(float) // // Given the transformed IR, a backend that supports the reg+immediate // addressing mode can easily fold the pointer arithmetics into the loads. For // example, the NVPTX backend can easily fold the pointer arithmetics into the // ld.global.f32 instructions, and the resultant PTX uses much fewer registers. // // mov.u32 %r1, %tid.x; // mov.u32 %r2, %tid.y; // mul.wide.u32 %rl2, %r1, 128; // mov.u64 %rl3, a; // add.s64 %rl4, %rl3, %rl2; // mul.wide.u32 %rl5, %r2, 4; // add.s64 %rl6, %rl4, %rl5; // ld.global.f32 %f1, [%rl6]; // so far the same as unoptimized PTX // ld.global.f32 %f2, [%rl6+4]; // much better // ld.global.f32 %f3, [%rl6+128]; // much better // ld.global.f32 %f4, [%rl6+132]; // much better // //===----------------------------------------------------------------------===// #include "llvm/Analysis/TargetTransformInfo.h" #include "llvm/Analysis/ValueTracking.h" #include "llvm/IR/Constants.h" #include "llvm/IR/DataLayout.h" #include "llvm/IR/Instructions.h" #include "llvm/IR/LLVMContext.h" #include "llvm/IR/Module.h" #include "llvm/IR/Operator.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/raw_ostream.h" #include "llvm/Transforms/Scalar.h" using namespace llvm; static cl::opt DisableSeparateConstOffsetFromGEP( "disable-separate-const-offset-from-gep", cl::init(false), cl::desc("Do not separate the constant offset from a GEP instruction"), cl::Hidden); namespace { /// \brief A helper class for separating a constant offset from a GEP index. /// /// In real programs, a GEP index may be more complicated than a simple addition /// of something and a constant integer which can be trivially splitted. For /// example, to split ((a << 3) | 5) + b, we need to search deeper for the /// constant offset, so that we can separate the index to (a << 3) + b and 5. /// /// Therefore, this class looks into the expression that computes a given GEP /// index, and tries to find a constant integer that can be hoisted to the /// outermost level of the expression as an addition. Not every constant in an /// expression can jump out. e.g., we cannot transform (b * (a + 5)) to (b * a + /// 5); nor can we transform (3 * (a + 5)) to (3 * a + 5), however in this case, /// -instcombine probably already optimized (3 * (a + 5)) to (3 * a + 15). class ConstantOffsetExtractor { public: /// Extracts a constant offset from the given GEP index. It outputs the /// numeric value of the extracted constant offset (0 if failed), and a /// new index representing the remainder (equal to the original index minus /// the constant offset). /// \p Idx The given GEP index /// \p NewIdx The new index to replace (output) /// \p DL The datalayout of the module /// \p GEP The given GEP static int64_t Extract(Value *Idx, Value *&NewIdx, const DataLayout *DL, GetElementPtrInst *GEP); /// Looks for a constant offset without extracting it. The meaning of the /// arguments and the return value are the same as Extract. static int64_t Find(Value *Idx, const DataLayout *DL, GetElementPtrInst *GEP); private: ConstantOffsetExtractor(const DataLayout *Layout, Instruction *InsertionPt) : DL(Layout), IP(InsertionPt) {} /// Searches the expression that computes V for a non-zero constant C s.t. /// V can be reassociated into the form V' + C. If the searching is /// successful, returns C and update UserChain as a def-use chain from C to V; /// otherwise, UserChain is empty. /// /// \p V The given expression /// \p SignExtended Whether V will be sign-extended in the computation of the /// GEP index /// \p ZeroExtended Whether V will be zero-extended in the computation of the /// GEP index /// \p NonNegative Whether V is guaranteed to be non-negative. For example, /// an index of an inbounds GEP is guaranteed to be /// non-negative. Levaraging this, we can better split /// inbounds GEPs. APInt find(Value *V, bool SignExtended, bool ZeroExtended, bool NonNegative); /// A helper function to look into both operands of a binary operator. APInt findInEitherOperand(BinaryOperator *BO, bool SignExtended, bool ZeroExtended); /// After finding the constant offset C from the GEP index I, we build a new /// index I' s.t. I' + C = I. This function builds and returns the new /// index I' according to UserChain produced by function "find". /// /// The building conceptually takes two steps: /// 1) iteratively distribute s/zext towards the leaves of the expression tree /// that computes I /// 2) reassociate the expression tree to the form I' + C. /// /// For example, to extract the 5 from sext(a + (b + 5)), we first distribute /// sext to a, b and 5 so that we have /// sext(a) + (sext(b) + 5). /// Then, we reassociate it to /// (sext(a) + sext(b)) + 5. /// Given this form, we know I' is sext(a) + sext(b). Value *rebuildWithoutConstOffset(); /// After the first step of rebuilding the GEP index without the constant /// offset, distribute s/zext to the operands of all operators in UserChain. /// e.g., zext(sext(a + (b + 5)) (assuming no overflow) => /// zext(sext(a)) + (zext(sext(b)) + zext(sext(5))). /// /// The function also updates UserChain to point to new subexpressions after /// distributing s/zext. e.g., the old UserChain of the above example is /// 5 -> b + 5 -> a + (b + 5) -> sext(...) -> zext(sext(...)), /// and the new UserChain is /// zext(sext(5)) -> zext(sext(b)) + zext(sext(5)) -> /// zext(sext(a)) + (zext(sext(b)) + zext(sext(5)) /// /// \p ChainIndex The index to UserChain. ChainIndex is initially /// UserChain.size() - 1, and is decremented during /// the recursion. Value *distributeExtsAndCloneChain(unsigned ChainIndex); /// Reassociates the GEP index to the form I' + C and returns I'. Value *removeConstOffset(unsigned ChainIndex); /// A helper function to apply ExtInsts, a list of s/zext, to value V. /// e.g., if ExtInsts = [sext i32 to i64, zext i16 to i32], this function /// returns "sext i32 (zext i16 V to i32) to i64". Value *applyExts(Value *V); /// Returns true if LHS and RHS have no bits in common, i.e., LHS | RHS == 0. bool NoCommonBits(Value *LHS, Value *RHS) const; /// Computes which bits are known to be one or zero. /// \p KnownOne Mask of all bits that are known to be one. /// \p KnownZero Mask of all bits that are known to be zero. void ComputeKnownBits(Value *V, APInt &KnownOne, APInt &KnownZero) const; /// A helper function that returns whether we can trace into the operands /// of binary operator BO for a constant offset. /// /// \p SignExtended Whether BO is surrounded by sext /// \p ZeroExtended Whether BO is surrounded by zext /// \p NonNegative Whether BO is known to be non-negative, e.g., an in-bound /// array index. bool CanTraceInto(bool SignExtended, bool ZeroExtended, BinaryOperator *BO, bool NonNegative); /// The path from the constant offset to the old GEP index. e.g., if the GEP /// index is "a * b + (c + 5)". After running function find, UserChain[0] will /// be the constant 5, UserChain[1] will be the subexpression "c + 5", and /// UserChain[2] will be the entire expression "a * b + (c + 5)". /// /// This path helps to rebuild the new GEP index. SmallVector UserChain; /// A data structure used in rebuildWithoutConstOffset. Contains all /// sext/zext instructions along UserChain. SmallVector ExtInsts; /// The data layout of the module. Used in ComputeKnownBits. const DataLayout *DL; Instruction *IP; /// Insertion position of cloned instructions. }; /// \brief A pass that tries to split every GEP in the function into a variadic /// base and a constant offset. It is a FunctionPass because searching for the /// constant offset may inspect other basic blocks. class SeparateConstOffsetFromGEP : public FunctionPass { public: static char ID; SeparateConstOffsetFromGEP() : FunctionPass(ID) { initializeSeparateConstOffsetFromGEPPass(*PassRegistry::getPassRegistry()); } void getAnalysisUsage(AnalysisUsage &AU) const override { 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: /// Tries to split the given GEP into a variadic base and a constant offset, /// and returns true if the splitting succeeds. bool splitGEP(GetElementPtrInst *GEP); /// Finds the constant offset within each index, and accumulates them. This /// 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, 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. /// /// Verified in @i32_add in split-gep.ll bool canonicalizeArrayIndicesToPointerSize(GetElementPtrInst *GEP); const DataLayout *DL; }; } // anonymous namespace char SeparateConstOffsetFromGEP::ID = 0; INITIALIZE_PASS_BEGIN( SeparateConstOffsetFromGEP, "separate-const-offset-from-gep", "Split GEPs to a variadic base and a constant offset for better CSE", false, false) INITIALIZE_AG_DEPENDENCY(TargetTransformInfo) INITIALIZE_PASS_DEPENDENCY(DataLayoutPass) INITIALIZE_PASS_END( SeparateConstOffsetFromGEP, "separate-const-offset-from-gep", "Split GEPs to a variadic base and a constant offset for better CSE", false, false) FunctionPass *llvm::createSeparateConstOffsetFromGEPPass() { return new SeparateConstOffsetFromGEP(); } bool ConstantOffsetExtractor::CanTraceInto(bool SignExtended, bool ZeroExtended, BinaryOperator *BO, bool NonNegative) { // We only consider ADD, SUB and OR, because a non-zero constant found in // expressions composed of these operations can be easily hoisted as a // constant offset by reassociation. if (BO->getOpcode() != Instruction::Add && BO->getOpcode() != Instruction::Sub && BO->getOpcode() != Instruction::Or) { return false; } Value *LHS = BO->getOperand(0), *RHS = BO->getOperand(1); // Do not trace into "or" unless it is equivalent to "add". If LHS and RHS // don't have common bits, (LHS | RHS) is equivalent to (LHS + RHS). if (BO->getOpcode() == Instruction::Or && !NoCommonBits(LHS, RHS)) return false; // In addition, tracing into BO requires that its surrounding s/zext (if // any) is distributable to both operands. // // Suppose BO = A op B. // SignExtended | ZeroExtended | Distributable? // --------------+--------------+---------------------------------- // 0 | 0 | true because no s/zext exists // 0 | 1 | zext(BO) == zext(A) op zext(B) // 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 && !ZeroExtended && NonNegative) { // If a + b >= 0 and (a >= 0 or b >= 0), then // 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 // index if the constant offset is non-negative. // // Verified in @sext_add in split-gep.ll. if (ConstantInt *ConstLHS = dyn_cast(LHS)) { if (!ConstLHS->isNegative()) return true; } if (ConstantInt *ConstRHS = dyn_cast(RHS)) { if (!ConstRHS->isNegative()) return true; } } // sext (add/sub nsw A, B) == add/sub nsw (sext A), (sext B) // zext (add/sub nuw A, B) == add/sub nuw (zext A), (zext B) if (BO->getOpcode() == Instruction::Add || BO->getOpcode() == Instruction::Sub) { if (SignExtended && !BO->hasNoSignedWrap()) return false; if (ZeroExtended && !BO->hasNoUnsignedWrap()) return false; } return true; } APInt ConstantOffsetExtractor::findInEitherOperand(BinaryOperator *BO, bool SignExtended, bool ZeroExtended) { // BO being non-negative does not shed light on whether its operands are // non-negative. Clear the NonNegative flag here. APInt ConstantOffset = find(BO->getOperand(0), SignExtended, ZeroExtended, /* NonNegative */ false); // If we found a constant offset in the left operand, stop and return that. // This shortcut might cause us to miss opportunities of combining the // constant offsets in both operands, e.g., (a + 4) + (b + 5) => (a + b) + 9. // However, such cases are probably already handled by -instcombine, // given this pass runs after the standard optimizations. if (ConstantOffset != 0) return ConstantOffset; ConstantOffset = find(BO->getOperand(1), SignExtended, ZeroExtended, /* NonNegative */ false); // If U is a sub operator, negate the constant offset found in the right // operand. if (BO->getOpcode() == Instruction::Sub) ConstantOffset = -ConstantOffset; return ConstantOffset; } APInt ConstantOffsetExtractor::find(Value *V, bool SignExtended, bool ZeroExtended, bool NonNegative) { // TODO(jingyue): We could trace into integer/pointer casts, such as // inttoptr, ptrtoint, bitcast, and addrspacecast. We choose to handle only // integers because it gives good enough results for our benchmarks. unsigned BitWidth = cast(V->getType())->getBitWidth(); // We cannot do much with Values that are not a User, such as an Argument. User *U = dyn_cast(V); if (U == nullptr) return APInt(BitWidth, 0); APInt ConstantOffset(BitWidth, 0); if (ConstantInt *CI = dyn_cast(V)) { // Hooray, we found it! ConstantOffset = CI->getValue(); } else if (BinaryOperator *BO = dyn_cast(V)) { // Trace into subexpressions for more hoisting opportunities. if (CanTraceInto(SignExtended, ZeroExtended, BO, NonNegative)) { ConstantOffset = findInEitherOperand(BO, SignExtended, ZeroExtended); } } else if (isa(V)) { ConstantOffset = find(U->getOperand(0), /* SignExtended */ true, ZeroExtended, NonNegative).sext(BitWidth); } else if (isa(V)) { // As an optimization, we can clear the SignExtended flag because // 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); } // If we found a non-zero constant offset, add it to the path for // rebuildWithoutConstOffset. Zero is a valid constant offset, but doesn't // help this optimization. if (ConstantOffset != 0) UserChain.push_back(U); return ConstantOffset; } Value *ConstantOffsetExtractor::applyExts(Value *V) { Value *Current = V; // ExtInsts is built in the use-def order. Therefore, we apply them to V // in the reversed order. for (auto I = ExtInsts.rbegin(), E = ExtInsts.rend(); I != E; ++I) { if (Constant *C = dyn_cast(Current)) { // If Current is a constant, apply s/zext using ConstantExpr::getCast. // ConstantExpr::getCast emits a ConstantInt if C is a ConstantInt. Current = ConstantExpr::getCast((*I)->getOpcode(), C, (*I)->getType()); } else { Instruction *Ext = (*I)->clone(); Ext->setOperand(0, Current); Ext->insertBefore(IP); Current = Ext; } } return Current; } Value *ConstantOffsetExtractor::rebuildWithoutConstOffset() { distributeExtsAndCloneChain(UserChain.size() - 1); // Remove all nullptrs (used to be s/zext) from UserChain. unsigned NewSize = 0; for (auto I = UserChain.begin(), E = UserChain.end(); I != E; ++I) { if (*I != nullptr) { UserChain[NewSize] = *I; NewSize++; } } UserChain.resize(NewSize); return removeConstOffset(UserChain.size() - 1); } Value * ConstantOffsetExtractor::distributeExtsAndCloneChain(unsigned ChainIndex) { User *U = UserChain[ChainIndex]; if (ChainIndex == 0) { assert(isa(U)); // If U is a ConstantInt, applyExts will return a ConstantInt as well. return UserChain[ChainIndex] = cast(applyExts(U)); } if (CastInst *Cast = dyn_cast(U)) { assert((isa(Cast) || isa(Cast)) && "We only traced into two types of CastInst: sext and zext"); ExtInsts.push_back(Cast); UserChain[ChainIndex] = nullptr; return distributeExtsAndCloneChain(ChainIndex - 1); } // Function find only trace into BinaryOperator and CastInst. BinaryOperator *BO = cast(U); // OpNo = which operand of BO is UserChain[ChainIndex - 1] unsigned OpNo = (BO->getOperand(0) == UserChain[ChainIndex - 1] ? 0 : 1); Value *TheOther = applyExts(BO->getOperand(1 - OpNo)); Value *NextInChain = distributeExtsAndCloneChain(ChainIndex - 1); BinaryOperator *NewBO = nullptr; if (OpNo == 0) { NewBO = BinaryOperator::Create(BO->getOpcode(), NextInChain, TheOther, BO->getName(), IP); } else { NewBO = BinaryOperator::Create(BO->getOpcode(), TheOther, NextInChain, BO->getName(), IP); } return UserChain[ChainIndex] = NewBO; } Value *ConstantOffsetExtractor::removeConstOffset(unsigned ChainIndex) { if (ChainIndex == 0) { assert(isa(UserChain[ChainIndex])); return ConstantInt::getNullValue(UserChain[ChainIndex]->getType()); } BinaryOperator *BO = cast(UserChain[ChainIndex]); unsigned OpNo = (BO->getOperand(0) == UserChain[ChainIndex - 1] ? 0 : 1); assert(BO->getOperand(OpNo) == UserChain[ChainIndex - 1]); Value *NextInChain = removeConstOffset(ChainIndex - 1); Value *TheOther = BO->getOperand(1 - OpNo); // If NextInChain is 0 and not the LHS of a sub, we can simplify the // sub-expression to be just TheOther. if (ConstantInt *CI = dyn_cast(NextInChain)) { if (CI->isZero() && !(BO->getOpcode() == Instruction::Sub && OpNo == 0)) return TheOther; } if (BO->getOpcode() == Instruction::Or) { // Rebuild "or" as "add", because "or" may be invalid for the new // epxression. // // For instance, given // a | (b + 5) where a and b + 5 have no common bits, // we can extract 5 as the constant offset. // // However, reusing the "or" in the new index would give us // (a | b) + 5 // which does not equal a | (b + 5). // // Replacing the "or" with "add" is fine, because // a | (b + 5) = a + (b + 5) = (a + b) + 5 return BinaryOperator::CreateAdd(BO->getOperand(0), BO->getOperand(1), BO->getName(), IP); } // We can reuse BO in this case, because the new expression shares the same // instruction type and BO is used at most once. assert(BO->getNumUses() <= 1 && "distributeExtsAndCloneChain clones each BinaryOperator in " "UserChain, so no one should be used more than " "once"); BO->setOperand(OpNo, NextInChain); BO->setHasNoSignedWrap(false); BO->setHasNoUnsignedWrap(false); // Make sure it appears after all instructions we've inserted so far. BO->moveBefore(IP); return BO; } int64_t ConstantOffsetExtractor::Extract(Value *Idx, Value *&NewIdx, const DataLayout *DL, GetElementPtrInst *GEP) { ConstantOffsetExtractor Extractor(DL, GEP); // Find a non-zero constant offset first. APInt ConstantOffset = Extractor.find(Idx, /* SignExtended */ false, /* ZeroExtended */ false, GEP->isInBounds()); if (ConstantOffset != 0) { // Separates the constant offset from the GEP index. NewIdx = Extractor.rebuildWithoutConstOffset(); } return ConstantOffset.getSExtValue(); } int64_t ConstantOffsetExtractor::Find(Value *Idx, const DataLayout *DL, GetElementPtrInst *GEP) { // If Idx is an index of an inbound GEP, Idx is guaranteed to be non-negative. return ConstantOffsetExtractor(DL, GEP) .find(Idx, /* SignExtended */ false, /* ZeroExtended */ false, GEP->isInBounds()) .getSExtValue(); } void ConstantOffsetExtractor::ComputeKnownBits(Value *V, APInt &KnownOne, APInt &KnownZero) const { IntegerType *IT = cast(V->getType()); KnownOne = APInt(IT->getBitWidth(), 0); KnownZero = APInt(IT->getBitWidth(), 0); llvm::computeKnownBits(V, KnownZero, KnownOne, DL, 0); } bool ConstantOffsetExtractor::NoCommonBits(Value *LHS, Value *RHS) const { assert(LHS->getType() == RHS->getType() && "LHS and RHS should have the same type"); APInt LHSKnownOne, LHSKnownZero, RHSKnownOne, RHSKnownZero; ComputeKnownBits(LHS, LHSKnownOne, LHSKnownZero); ComputeKnownBits(RHS, RHSKnownOne, RHSKnownZero); return (LHSKnownZero | RHSKnownZero).isAllOnesValue(); } 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; } int64_t SeparateConstOffsetFromGEP::accumulateByteOffset(GetElementPtrInst *GEP, bool &NeedsExtraction) { NeedsExtraction = false; int64_t AccumulativeByteOffset = 0; gep_type_iterator GTI = gep_type_begin(*GEP); for (unsigned I = 1, E = GEP->getNumOperands(); I != E; ++I, ++GTI) { if (isa(*GTI)) { // Tries to extract a constant offset from this GEP index. int64_t ConstantOffset = ConstantOffsetExtractor::Find(GEP->getOperand(I), DL, GEP); if (ConstantOffset != 0) { NeedsExtraction = true; // A GEP may have multiple indices. We accumulate the extracted // constant offset to a byte offset, and later offset the remainder of // the original GEP with this byte offset. AccumulativeByteOffset += ConstantOffset * DL->getTypeAllocSize(GTI.getIndexedType()); } } } return AccumulativeByteOffset; } bool SeparateConstOffsetFromGEP::splitGEP(GetElementPtrInst *GEP) { // Skip vector GEPs. if (GEP->getType()->isVectorTy()) return false; // The backend can already nicely handle the case where all indices are // constant. if (GEP->hasAllConstantIndices()) return false; bool Changed = false; Changed |= canonicalizeArrayIndicesToPointerSize(GEP); bool NeedsExtraction; int64_t AccumulativeByteOffset = accumulateByteOffset(GEP, NeedsExtraction); if (!NeedsExtraction) return Changed; // Before really splitting the GEP, check whether the backend supports the // addressing mode we are about to produce. If no, this splitting probably // won't be beneficial. TargetTransformInfo &TTI = getAnalysis(); if (!TTI.isLegalAddressingMode(GEP->getType()->getElementType(), /*BaseGV=*/nullptr, AccumulativeByteOffset, /*HasBaseReg=*/true, /*Scale=*/0)) { return Changed; } // Remove the constant offset in each GEP index. The resultant GEP computes // the variadic base. 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; // Tries to extract a constant offset from this GEP index. int64_t ConstantOffset = ConstantOffsetExtractor::Extract(GEP->getOperand(I), NewIdx, DL, GEP); if (ConstantOffset != 0) { assert(NewIdx != nullptr && "ConstantOffset != 0 implies NewIdx is set"); GEP->setOperand(I, NewIdx); } } } // Clear the inbounds attribute because the new index may be off-bound. // e.g., // // b = add i64 a, 5 // addr = gep inbounds float* p, i64 b // // is transformed to: // // addr2 = gep float* p, i64 a // addr = gep float* addr2, i64 5 // // If a is -4, although the old index b is in bounds, the new index a is // off-bound. http://llvm.org/docs/LangRef.html#id181 says "if the // inbounds keyword is not present, the offsets are added to the base // address with silently-wrapping two's complement arithmetic". // Therefore, the final code will be a semantically equivalent. // // TODO(jingyue): do some range analysis to keep as many inbounds as // possible. GEPs with inbounds are more friendly to alias analysis. GEP->setIsInBounds(false); // Offsets the base with the accumulative byte offset. // // %gep ; the base // ... %gep ... // // => add the offset // // %gep2 ; clone of %gep // %new.gep = gep %gep2, // %gep ; will be removed // ... %gep ... // // => replace all uses of %gep with %new.gep and remove %gep // // %gep2 ; clone of %gep // %new.gep = gep %gep2, // ... %new.gep ... // // If AccumulativeByteOffset is not a multiple of sizeof(*%gep), we emit an // uglygep (http://llvm.org/docs/GetElementPtr.html#what-s-an-uglygep): // bitcast %gep2 to i8*, add the offset, and bitcast the result back to the // type of %gep. // // %gep2 ; clone of %gep // %0 = bitcast %gep2 to i8* // %uglygep = gep %0, // %new.gep = bitcast %uglygep to // ... %new.gep ... Instruction *NewGEP = GEP->clone(); NewGEP->insertBefore(GEP); 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). // Per ANSI C standard, signed / unsigned = unsigned. Therefore, we // cast ElementTypeSizeOfGEP to signed. int64_t Index = AccumulativeByteOffset / static_cast(ElementTypeSizeOfGEP); NewGEP = GetElementPtrInst::Create( NewGEP, ConstantInt::get(IntPtrTy, Index, true), GEP->getName(), GEP); } else { // Unlikely but possible. For example, // #pragma pack(1) // struct S { // int a[3]; // int64 b[8]; // }; // #pragma pack() // // Suppose the gep before extraction is &s[i + 1].b[j + 3]. After // extraction, it becomes &s[i].b[j] and AccumulativeByteOffset is // sizeof(S) + 3 * sizeof(int64) = 100, which is not a multiple of // sizeof(int64). // // Emit an uglygep in this case. Type *I8PtrTy = Type::getInt8PtrTy(GEP->getContext(), GEP->getPointerAddressSpace()); NewGEP = new BitCastInst(NewGEP, I8PtrTy, "", GEP); NewGEP = GetElementPtrInst::Create( NewGEP, ConstantInt::get(IntPtrTy, AccumulativeByteOffset, true), "uglygep", GEP); if (GEP->getType() != I8PtrTy) NewGEP = new BitCastInst(NewGEP, GEP->getType(), GEP->getName(), GEP); } GEP->replaceAllUsesWith(NewGEP); GEP->eraseFromParent(); return true; } bool SeparateConstOffsetFromGEP::runOnFunction(Function &F) { if (DisableSeparateConstOffsetFromGEP) return false; bool Changed = false; for (Function::iterator B = F.begin(), BE = F.end(); B != BE; ++B) { for (BasicBlock::iterator I = B->begin(), IE = B->end(); I != IE; ) { if (GetElementPtrInst *GEP = dyn_cast(I++)) { Changed |= splitGEP(GEP); } // No need to split GEP ConstantExprs because all its indices are constant // already. } } return Changed; }