1//===- SeparateConstOffsetFromGEP.cpp -------------------------------------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9// Loop unrolling may create many similar GEPs for array accesses.
10// e.g., a 2-level loop
11//
12// float a[32][32]; // global variable
13//
14// for (int i = 0; i < 2; ++i) {
15// for (int j = 0; j < 2; ++j) {
16// ...
17// ... = a[x + i][y + j];
18// ...
19// }
20// }
21//
22// will probably be unrolled to:
23//
24// gep %a, 0, %x, %y; load
25// gep %a, 0, %x, %y + 1; load
26// gep %a, 0, %x + 1, %y; load
27// gep %a, 0, %x + 1, %y + 1; load
28//
29// LLVM's GVN does not use partial redundancy elimination yet, and is thus
30// unable to reuse (gep %a, 0, %x, %y). As a result, this misoptimization incurs
31// significant slowdown in targets with limited addressing modes. For instance,
32// because the PTX target does not support the reg+reg addressing mode, the
33// NVPTX backend emits PTX code that literally computes the pointer address of
34// each GEP, wasting tons of registers. It emits the following PTX for the
35// first load and similar PTX for other loads.
36//
37// mov.u32 %r1, %x;
38// mov.u32 %r2, %y;
39// mul.wide.u32 %rl2, %r1, 128;
40// mov.u64 %rl3, a;
41// add.s64 %rl4, %rl3, %rl2;
42// mul.wide.u32 %rl5, %r2, 4;
43// add.s64 %rl6, %rl4, %rl5;
44// ld.global.f32 %f1, [%rl6];
45//
46// To reduce the register pressure, the optimization implemented in this file
47// merges the common part of a group of GEPs, so we can compute each pointer
48// address by adding a simple offset to the common part, saving many registers.
49//
50// It works by splitting each GEP into a variadic base and a constant offset.
51// The variadic base can be computed once and reused by multiple GEPs, and the
52// constant offsets can be nicely folded into the reg+immediate addressing mode
53// (supported by most targets) without using any extra register.
54//
55// For instance, we transform the four GEPs and four loads in the above example
56// into:
57//
58// base = gep a, 0, x, y
59// load base
60// load base + 1 * sizeof(float)
61// load base + 32 * sizeof(float)
62// load base + 33 * sizeof(float)
63//
64// Given the transformed IR, a backend that supports the reg+immediate
65// addressing mode can easily fold the pointer arithmetics into the loads. For
66// example, the NVPTX backend can easily fold the pointer arithmetics into the
67// ld.global.f32 instructions, and the resultant PTX uses much fewer registers.
68//
69// mov.u32 %r1, %tid.x;
70// mov.u32 %r2, %tid.y;
71// mul.wide.u32 %rl2, %r1, 128;
72// mov.u64 %rl3, a;
73// add.s64 %rl4, %rl3, %rl2;
74// mul.wide.u32 %rl5, %r2, 4;
75// add.s64 %rl6, %rl4, %rl5;
76// ld.global.f32 %f1, [%rl6]; // so far the same as unoptimized PTX
77// ld.global.f32 %f2, [%rl6+4]; // much better
78// ld.global.f32 %f3, [%rl6+128]; // much better
79// ld.global.f32 %f4, [%rl6+132]; // much better
80//
81// Another improvement enabled by the LowerGEP flag is to lower a GEP with
82// multiple indices to multiple GEPs with a single index.
83// Such transformation can have following benefits:
84// (1) It can always extract constants in the indices of structure type.
85// (2) After such Lowering, there are more optimization opportunities such as
86// CSE, LICM and CGP.
87//
88// E.g. The following GEPs have multiple indices:
89// BB1:
90// %p = getelementptr [10 x %struct], ptr %ptr, i64 %i, i64 %j1, i32 3
91// load %p
92// ...
93// BB2:
94// %p2 = getelementptr [10 x %struct], ptr %ptr, i64 %i, i64 %j1, i32 2
95// load %p2
96// ...
97//
98// We can not do CSE to the common part related to index "i64 %i". Lowering
99// GEPs can achieve such goals.
100//
101// This pass will lower a GEP with multiple indices into multiple GEPs with a
102// single index:
103// BB1:
104// %2 = mul i64 %i, length_of_10xstruct ; CSE opportunity
105// %3 = getelementptr i8, ptr %ptr, i64 %2 ; CSE opportunity
106// %4 = mul i64 %j1, length_of_struct
107// %5 = getelementptr i8, ptr %3, i64 %4
108// %p = getelementptr i8, ptr %5, struct_field_3 ; Constant offset
109// load %p
110// ...
111// BB2:
112// %8 = mul i64 %i, length_of_10xstruct ; CSE opportunity
113// %9 = getelementptr i8, ptr %ptr, i64 %8 ; CSE opportunity
114// %10 = mul i64 %j2, length_of_struct
115// %11 = getelementptr i8, ptr %9, i64 %10
116// %p2 = getelementptr i8, ptr %11, struct_field_2 ; Constant offset
117// load %p2
118// ...
119//
120// Lowering GEPs can also benefit other passes such as LICM and CGP.
121// LICM (Loop Invariant Code Motion) can not hoist/sink a GEP of multiple
122// indices if one of the index is variant. If we lower such GEP into invariant
123// parts and variant parts, LICM can hoist/sink those invariant parts.
124// CGP (CodeGen Prepare) tries to sink address calculations that match the
125// target's addressing modes. A GEP with multiple indices may not match and will
126// not be sunk. If we lower such GEP into smaller parts, CGP may sink some of
127// them. So we end up with a better addressing mode.
128//
129//===----------------------------------------------------------------------===//
130
131#include "llvm/Transforms/Scalar/SeparateConstOffsetFromGEP.h"
132#include "ScalarOptions.h"
133#include "llvm/ADT/APInt.h"
134#include "llvm/ADT/DenseMap.h"
135#include "llvm/ADT/DepthFirstIterator.h"
136#include "llvm/ADT/SmallVector.h"
137#include "llvm/Analysis/LoopInfo.h"
138#include "llvm/Analysis/MemoryBuiltins.h"
139#include "llvm/Analysis/TargetLibraryInfo.h"
140#include "llvm/Analysis/TargetTransformInfo.h"
141#include "llvm/Analysis/ValueTracking.h"
142#include "llvm/IR/BasicBlock.h"
143#include "llvm/IR/Constant.h"
144#include "llvm/IR/Constants.h"
145#include "llvm/IR/DataLayout.h"
146#include "llvm/IR/DerivedTypes.h"
147#include "llvm/IR/Dominators.h"
148#include "llvm/IR/Function.h"
149#include "llvm/IR/GetElementPtrTypeIterator.h"
150#include "llvm/IR/IRBuilder.h"
151#include "llvm/IR/InstrTypes.h"
152#include "llvm/IR/Instruction.h"
153#include "llvm/IR/Instructions.h"
154#include "llvm/IR/Module.h"
155#include "llvm/IR/PassManager.h"
156#include "llvm/IR/PatternMatch.h"
157#include "llvm/IR/Type.h"
158#include "llvm/IR/User.h"
159#include "llvm/IR/Value.h"
160#include "llvm/InitializePasses.h"
161#include "llvm/Pass.h"
162#include "llvm/Support/Casting.h"
163#include "llvm/Support/ErrorHandling.h"
164#include "llvm/Support/KnownBits.h"
165#include "llvm/Support/raw_ostream.h"
166#include "llvm/Transforms/Scalar.h"
167#include "llvm/Transforms/Utils/Local.h"
168#include <cassert>
169#include <cstdint>
170#include <iterator>
171#include <optional>
172#include <string>
173
174using namespace llvm;
175using namespace llvm::PatternMatch;
176
177namespace {
178
179/// A helper class for separating a constant offset from a GEP index.
180///
181/// In real programs, a GEP index may be more complicated than a simple addition
182/// of something and a constant integer which can be trivially splitted. For
183/// example, to split ((a << 3) | 5) + b, we need to search deeper for the
184/// constant offset, so that we can separate the index to (a << 3) + b and 5.
185///
186/// Therefore, this class looks into the expression that computes a given GEP
187/// index, and tries to find a constant integer that can be hoisted to the
188/// outermost level of the expression as an addition. Not every constant in an
189/// expression can jump out. e.g., we cannot transform (b * (a + 5)) to (b * a +
190/// 5); nor can we transform (3 * (a + 5)) to (3 * a + 5), however in this case,
191/// -instcombine probably already optimized (3 * (a + 5)) to (3 * a + 15).
192class ConstantOffsetExtractor {
193public:
194 /// Extracts a constant offset from the given GEP index. It returns the
195 /// new index representing the remainder (equal to the original index minus
196 /// the constant offset), or nullptr if we cannot extract a constant offset.
197 /// \p Idx The given GEP index use
198 /// \p UserTreeRoot Outputs the rebuilt index for dead instruction cleanup.
199 /// \p PreservesNUW Outputs whether the extraction allows preserving the
200 /// GEP's nuw flag, if it has one.
201 static Value *Extract(const Use &Idx, User *&UserTreeRoot,
202 bool &PreservesNUW);
203
204 /// Looks for a constant offset from the given GEP index without extracting
205 /// it. It returns the numeric value of the extracted constant offset, or
206 /// std::nullopt on failure. The arguments have the same meaning as Extract.
207 static std::optional<APInt> Find(const Use &Idx);
208
209private:
210 struct RewriteNode {
211 User *V;
212 std::optional<unsigned> LHS;
213 std::optional<unsigned> RHS;
214 };
215
216 struct CastState {
217 SmallVector<CastInst *, 4> Casts;
218
219 void pushCast(CastInst *Cast) {
220 assert((isa<SExtInst, ZExtInst, TruncInst>(Cast)) && "Unexpected cast");
221 Casts.push_back(Elt: Cast);
222 }
223
224 void popCast() { Casts.pop_back(); }
225
226 bool hasSignExtension() const {
227 // The innermost extension determines whether signed overflow matters:
228 // sext(zext(a)) = zext(a).
229 for (CastInst *C : reverse(C: Casts))
230 if (isa<SExtInst, ZExtInst>(Val: C))
231 return isa<SExtInst>(Val: C);
232 return false;
233 }
234
235 bool hasZeroExtension() const { return any_of(Range: Casts, P: IsaPred<ZExtInst>); }
236
237 bool hasExtensions() const {
238 return any_of(Range: Casts, P: IsaPred<SExtInst, ZExtInst>);
239 }
240
241 APInt apply(APInt Offset) const {
242 for (CastInst *Cast : llvm::reverse(C: Casts)) {
243 unsigned BitWidth = cast<IntegerType>(Val: Cast->getType())->getBitWidth();
244 if (isa<SExtInst>(Val: Cast))
245 Offset = Offset.sext(width: BitWidth);
246 else if (isa<ZExtInst>(Val: Cast))
247 Offset = Offset.zext(width: BitWidth);
248 else {
249 assert(isa<TruncInst>(Cast) && "Unexpected cast");
250 Offset = Offset.trunc(width: BitWidth);
251 }
252 }
253 return Offset;
254 }
255 };
256
257 ConstantOffsetExtractor(BasicBlock::iterator InsertionPt)
258 : IP(InsertionPt), DL(InsertionPt->getDataLayout()), SQ(DL) {}
259
260 /// Searches the expression that computes V for a constant offset C s.t.
261 /// V can be reassociated into the form V' + C. If the searching is
262 /// successful, returns C and appends the rewrite for V to UserTree;
263 /// otherwise, returns std::nullopt and leaves UserTree unchanged.
264 /// \p V The given expression
265 /// \p Idx The original index use of the GEP, or nullptr if its
266 /// sign and bounds information no longer applies
267 /// \p Casts The casts surrounding V in the original expression.
268 std::optional<APInt> find(Value *V, const Use *Idx, CastState &Casts);
269
270 /// Collect offsets from the operands and record their replacements.
271 std::optional<APInt> findInOperands(RewriteNode &Node, CastState &Casts);
272
273 /// After finding the constant offset C from the GEP index I, we build a new
274 /// index I' s.t. I' + C = I. This function builds and returns the new
275 /// index I' according to UserTree produced by function "find".
276 ///
277 /// While rebuilding, distribute the casts recorded in \p Casts to each
278 /// operand that remains in the expression, then reassociate the expression to
279 /// the form I' + C and return I'.
280 Value *rebuildWithoutConstOffset(unsigned NodeIndex, CastState &Casts);
281
282 Value *rebuildWithoutConstOffset() {
283 CastState Casts;
284 return rebuildWithoutConstOffset(NodeIndex: UserTree.size() - 1, Casts);
285 }
286
287 /// A helper function to apply a list of sext/zext/trunc casts to value
288 /// V. e.g., if Casts = [sext i32 to i64, zext i16 to i32], this function
289 /// returns "sext i32 (zext i16 V to i32) to i64".
290 Value *applyCasts(Value *V, const CastState &Casts);
291
292 /// A helper function that returns whether we can trace into the operands
293 /// of binary operator BO for a constant offset.
294 ///
295 /// \p Casts The casts surrounding BO.
296 /// \p Idx The original index use of the GEP, or nullptr if its
297 /// sign and bounds information no longer applies
298 bool canTraceInto(const CastState &Casts, BinaryOperator *BO, const Use *Idx);
299
300 /// Analyze a xor expression, and identify the bits in the constant operand
301 /// that are disjoint from the base operand's known set bits. For these
302 /// disjoint bits, a xor is equivalent to an addition, which allows us to
303 /// extract them as constant offsets that can be folded into the immediate
304 /// field of addressing operations. The transformation is the following one:
305 ///
306 /// Base ^ Const becomes (Base ^ NonDisjointBits) + DisjointBits
307 ///
308 /// where DisjointBits = Const & KnownZeros(Base) and
309 /// NonDisjointBits = Const & ~DisjointBits.
310 ///
311 /// Example with ptr having known-zero low bit:
312 /// Original: `xor %ptr, 3` ; 3 = 0b11
313 /// Analysis: DisjointBits = 3 & KnownZeros(%ptr) = 0b11 & 0b01 = 0b01
314 /// Result: `(xor %ptr, 2) + 1` where 1 can be folded into address mode
315 ///
316 /// \param Node The XOR node to analyze.
317 /// \return The disjoint bits (the extractable offset), or std::nullopt if
318 /// none exist. On success, records the remaining constant as a child node.
319 std::optional<APInt> extractDisjointBitsFromXor(RewriteNode &Node);
320
321 /// A postorder rewrite plan. Child indices refer to earlier entries; absent
322 /// children leave the original operands unchanged. Constant leaves hold
323 /// replacement values rather than the constants being extracted.
324 SmallVector<RewriteNode, 8> UserTree;
325
326 static constexpr unsigned MaxVisitedNodes = 64;
327
328 /// Count visits, including repeated visits to shared expressions.
329 unsigned NumVisited = 0;
330
331 /// Insertion position of cloned instructions.
332 BasicBlock::iterator IP;
333
334 const DataLayout &DL;
335 const SimplifyQuery SQ;
336};
337
338/// A pass that tries to split every GEP in the function into a variadic
339/// base and a constant offset. It is a FunctionPass because searching for the
340/// constant offset may inspect other basic blocks.
341class SeparateConstOffsetFromGEPLegacyPass : public FunctionPass {
342public:
343 static char ID;
344
345 SeparateConstOffsetFromGEPLegacyPass(bool LowerGEP = false)
346 : FunctionPass(ID), LowerGEP(LowerGEP) {
347 initializeSeparateConstOffsetFromGEPLegacyPassPass(
348 *PassRegistry::getPassRegistry());
349 }
350
351 void getAnalysisUsage(AnalysisUsage &AU) const override {
352 AU.addRequired<DominatorTreeWrapperPass>();
353 AU.addRequired<TargetTransformInfoWrapperPass>();
354 AU.addRequired<LoopInfoWrapperPass>();
355 AU.setPreservesCFG();
356 AU.addRequired<TargetLibraryInfoWrapperPass>();
357 }
358
359 bool runOnFunction(Function &F) override;
360
361private:
362 bool LowerGEP;
363};
364
365/// A pass that tries to split every GEP in the function into a variadic
366/// base and a constant offset. It is a FunctionPass because searching for the
367/// constant offset may inspect other basic blocks.
368class SeparateConstOffsetFromGEP {
369public:
370 SeparateConstOffsetFromGEP(
371 DominatorTree *DT, LoopInfo *LI, TargetLibraryInfo *TLI,
372 function_ref<TargetTransformInfo &(Function &)> GetTTI, bool LowerGEP)
373 : DT(DT), LI(LI), TLI(TLI), GetTTI(GetTTI), LowerGEP(LowerGEP) {}
374
375 bool run(Function &F);
376
377private:
378 /// Track the operands of an add or sub.
379 using ExprKey = std::pair<Value *, Value *>;
380
381 /// Create a pair for use as a map key for a commutable operation.
382 static ExprKey createNormalizedCommutablePair(Value *A, Value *B) {
383 if (A < B)
384 return {A, B};
385 return {B, A};
386 }
387
388 /// Tries to split the given GEP into a variadic base and a constant offset,
389 /// and returns true if the splitting succeeds.
390 bool splitGEP(GetElementPtrInst *GEP);
391
392 /// Tries to reorder the given GEP with the GEP that produces the base if
393 /// doing so results in producing a constant offset as the outermost
394 /// index.
395 bool reorderGEP(GetElementPtrInst *GEP, TargetTransformInfo &TTI);
396
397 /// Lower a GEP with multiple indices into multiple GEPs with a single index.
398 /// Function splitGEP already split the original GEP into a variadic part and
399 /// a constant offset (i.e., AccumulativeByteOffset). This function lowers the
400 /// variadic part into a set of GEPs with a single index and applies
401 /// AccumulativeByteOffset to it.
402 /// \p Variadic The variadic part of the original GEP.
403 /// \p AccumulativeByteOffset The constant offset.
404 void lowerToSingleIndexGEPs(GetElementPtrInst *Variadic,
405 const APInt &AccumulativeByteOffset);
406
407 /// Finds the constant offset within each index and accumulates them. If
408 /// LowerGEP is true, it finds in indices of both sequential and structure
409 /// types, otherwise it only finds in sequential indices. The output
410 /// NeedsExtraction indicates whether we successfully find a constant
411 /// offset, and SignedOverflow indicates if calculating or removing offsets
412 /// may have signed overflow.
413 APInt accumulateByteOffset(GetElementPtrInst *GEP, bool &NeedsExtraction,
414 bool &SignedOverflow);
415
416 /// Canonicalize array indices to pointer-size integers. This helps to
417 /// simplify the logic of splitting a GEP. For example, if a + b is a
418 /// pointer-size integer, we have
419 /// gep base, a + b = gep (gep base, a), b
420 /// However, this equality may not hold if the size of a + b is smaller than
421 /// the pointer size, because LLVM conceptually sign-extends GEP indices to
422 /// pointer size before computing the address
423 /// (http://llvm.org/docs/LangRef.html#id181).
424 ///
425 /// This canonicalization is very likely already done in clang and
426 /// instcombine. Therefore, the program will probably remain the same.
427 ///
428 /// Returns true if the module changes.
429 ///
430 /// Verified in @i32_add in split-gep.ll
431 bool canonicalizeArrayIndicesToIndexSize(GetElementPtrInst *GEP);
432
433 /// Optimize sext(a)+sext(b) to sext(a+b) when a+b can't sign overflow.
434 /// SeparateConstOffsetFromGEP distributes a sext to leaves before extracting
435 /// the constant offset. After extraction, it becomes desirable to reunion the
436 /// distributed sexts. For example,
437 ///
438 /// &a[sext(i +nsw (j +nsw 5)]
439 /// => distribute &a[sext(i) +nsw (sext(j) +nsw 5)]
440 /// => constant extraction &a[sext(i) + sext(j)] + 5
441 /// => reunion &a[sext(i +nsw j)] + 5
442 bool reuniteExts(Function &F);
443
444 /// A helper that reunites sexts in an instruction.
445 bool reuniteExts(Instruction *I);
446
447 /// Find the closest dominator of <Dominatee> that is equivalent to <Key>.
448 Instruction *findClosestMatchingDominator(
449 ExprKey Key, Instruction *Dominatee,
450 DenseMap<ExprKey, SmallVector<Instruction *, 2>> &DominatingExprs);
451
452 /// Verify F is free of dead code.
453 void verifyNoDeadCode(Function &F);
454
455 bool hasMoreThanOneUseInLoop(Value *v, Loop *L);
456
457 // Swap the index operand of two GEP.
458 void swapGEPOperand(GetElementPtrInst *First, GetElementPtrInst *Second);
459
460 // Check if it is safe to swap operand of two GEP.
461 bool isLegalToSwapOperand(GetElementPtrInst *First, GetElementPtrInst *Second,
462 Loop *CurLoop);
463
464 const DataLayout *DL = nullptr;
465 DominatorTree *DT = nullptr;
466 LoopInfo *LI;
467 TargetLibraryInfo *TLI;
468 // Retrieved lazily since not always used.
469 function_ref<TargetTransformInfo &(Function &)> GetTTI;
470
471 /// Whether to lower a GEP with multiple indices into arithmetic operations or
472 /// multiple GEPs with a single index.
473 bool LowerGEP;
474
475 DenseMap<ExprKey, SmallVector<Instruction *, 2>> DominatingAdds;
476 DenseMap<ExprKey, SmallVector<Instruction *, 2>> DominatingSubs;
477};
478
479} // end anonymous namespace
480
481char SeparateConstOffsetFromGEPLegacyPass::ID = 0;
482
483INITIALIZE_PASS_BEGIN(
484 SeparateConstOffsetFromGEPLegacyPass, "separate-const-offset-from-gep",
485 "Split GEPs to a variadic base and a constant offset for better CSE", false,
486 false)
487INITIALIZE_PASS_DEPENDENCY(DominatorTreeWrapperPass)
488INITIALIZE_PASS_DEPENDENCY(ScalarEvolutionWrapperPass)
489INITIALIZE_PASS_DEPENDENCY(TargetTransformInfoWrapperPass)
490INITIALIZE_PASS_DEPENDENCY(LoopInfoWrapperPass)
491INITIALIZE_PASS_DEPENDENCY(TargetLibraryInfoWrapperPass)
492INITIALIZE_PASS_END(
493 SeparateConstOffsetFromGEPLegacyPass, "separate-const-offset-from-gep",
494 "Split GEPs to a variadic base and a constant offset for better CSE", false,
495 false)
496
497FunctionPass *llvm::createSeparateConstOffsetFromGEPPass(bool LowerGEP) {
498 return new SeparateConstOffsetFromGEPLegacyPass(LowerGEP);
499}
500
501// Checks if it is safe to reorder an add/sext result used in a GEP.
502//
503// An inbounds GEP does not guarantee that the index is non-negative.
504// This helper first checks whether value tracking proves that the add cannot
505// have signed overflow. If so, the transform is safe.
506// Second, it checks whether the GEP is inbounds and directly based on a global
507// or an alloca, which are required to prove futher transform validity.
508// If the GEP:
509// - Has a zero offset from the base and Idx is its first index, the index is
510// non-negative (any negative value would produce poison/UB)
511// - Has ObjectSize < (2^(N-1) - C + 1) * stride, where C is a constant from the
512// add, stride is the element size of Idx, and N is bitwidth of the add.
513// This is because with this pattern:
514// %add = add iN %val, C
515// %sext = sext iN %add to i64
516// %gep = getelementptr inbounds TYPE, %sext
517// The worst-case is when %val sign-flips to produce the smallest magnitude
518// negative value, at 2^(N-1)-1. In this case, the add/sext is -(2^(N-1)-C+1),
519// and the sext/add is 2^(N-1)+C-1 (2^N difference). The original add/sext
520// only produces a defined GEP when -(2^(N-1)-C+1) is inbounds. So, if
521// ObjectSize < (2^(N-1) - C + 1) * stride, it is impossible for the
522// worst-case sign-flip to be defined.
523// Note that in this case the GEP is not neccesarily non-negative, but any
524// negative results will still produce the same behavior in the reordered
525// version with a defined GEP.
526// This can also work for negative C, but the threshold is instead
527// (2^(N-1)+C)*stride, since the sign-flip is done in reverse and is instead
528// producing a large positive value that still needs to be inbounds to the
529// object size. If C is negative, we cannot make any useful assumptions based
530// on the offset, since it would need to be extremely large.
531static bool canReorderAddSextToGEP(const Use *Idx, const BinaryOperator *Add,
532 const DataLayout &DL) {
533 if (computeOverflowForSignedAdd(Add: cast<AddOperator>(Val: Add), SQ: DL) ==
534 OverflowResult::NeverOverflows)
535 return true;
536
537 const auto *GEP = cast<GetElementPtrInst>(Val: Idx->getUser());
538 if (!GEP->isInBounds())
539 return false;
540
541 const Value *Ptr = GEP->getPointerOperand();
542 int64_t Offset = 0;
543 const Value *Base =
544 GetPointerBaseWithConstantOffset(Ptr: const_cast<Value *>(Ptr), Offset, DL);
545
546 // We need one of the operands to be a constant to be able to trace into the
547 // operator.
548 const ConstantInt *CI = dyn_cast<ConstantInt>(Val: Add->getOperand(i_nocapture: 0));
549 if (!CI)
550 CI = dyn_cast<ConstantInt>(Val: Add->getOperand(i_nocapture: 1));
551 if (!CI)
552 return false;
553 // Calculate the threshold
554 APInt Threshold;
555 unsigned N = Add->getType()->getIntegerBitWidth();
556 // Track the use: the same value may index different types in this GEP.
557 auto GTI = std::next(x: gep_type_begin(GEP), n: Idx->getOperandNo() - 1);
558 TypeSize ElemSize = GTI.getSequentialElementStride(DL);
559 if (ElemSize.isScalable())
560 return false;
561 uint64_t Stride = ElemSize.getFixedValue();
562 if (!CI->isNegative()) {
563 // (2^(N-1) - C + 1) * stride
564 Threshold = (APInt::getSignedMinValue(numBits: N).zext(width: 128) -
565 CI->getValue().zextOrTrunc(width: 128) + 1) *
566 APInt(128, Stride);
567 } else {
568 // (2^(N-1) + C) * stride
569 Threshold = (APInt::getSignedMinValue(numBits: N).zext(width: 128) +
570 CI->getValue().sextOrTrunc(width: 128)) *
571 APInt(128, Stride);
572 }
573
574 // Only the first index is relative to Ptr. Earlier indices may move the
575 // pointer within the object, so later indices must use the object-size proof.
576 if (Idx->getOperandNo() == 1 && Base &&
577 (isa<AllocaInst>(Val: Base) || isa<GlobalObject>(Val: Base)) && !CI->isNegative()) {
578 // If the offset is zero from an alloca or global, inbounds is sufficient to
579 // prove non-negativity if one add operand is non-negative
580 if (Offset == 0)
581 return true;
582
583 // Check if the Offset < Threshold (positive CI only) otherwise
584 if (Offset < 0)
585 return true;
586 if (APInt(128, (uint64_t)Offset).ult(RHS: Threshold))
587 return true;
588 } else {
589 // If we can't determine the offset from the base object, we can still use
590 // the underlying object and type size constraints
591 Base = getUnderlyingObject(V: Ptr);
592 // Can only prove non-negativity if the base object is known
593 if (!(isa<AllocaInst>(Val: Base) || isa<GlobalObject>(Val: Base)))
594 return false;
595 }
596
597 // Check if the ObjectSize < Threshold (for both positive or negative C)
598 uint64_t ObjSize = 0;
599 if (const auto *AI = dyn_cast<AllocaInst>(Val: Base)) {
600 if (auto AllocSize = AI->getAllocationSize(DL))
601 if (!AllocSize->isScalable())
602 ObjSize = AllocSize->getFixedValue();
603 } else if (const auto *GV = dyn_cast<GlobalVariable>(Val: Base)) {
604 TypeSize GVSize = DL.getTypeAllocSize(Ty: GV->getValueType());
605 if (!GVSize.isScalable())
606 ObjSize = GVSize.getFixedValue();
607 }
608 if (ObjSize > 0 && APInt(128, ObjSize).ult(RHS: Threshold))
609 return true;
610
611 return false;
612}
613
614bool ConstantOffsetExtractor::canTraceInto(const CastState &Casts,
615 BinaryOperator *BO, const Use *Idx) {
616 bool SignExtended = Casts.hasSignExtension();
617 bool ZeroExtended = Casts.hasZeroExtension();
618 // Do not trace into "or" unless it is equivalent to "add nuw nsw".
619 // This is the case if the or's disjoint flag is set.
620 if (BO->getOpcode() == Instruction::Or)
621 return cast<PossiblyDisjointInst>(Val: BO)->isDisjoint();
622
623 // We only consider ADD and SUB here, because a non-zero constant found in
624 // expressions composed of these operations can be easily hoisted as a
625 // constant offset by reassociation.
626 if (BO->getOpcode() != Instruction::Add &&
627 BO->getOpcode() != Instruction::Sub)
628 return false;
629
630 // In addition, tracing into BO requires that its surrounding sext/zext/trunc
631 // (if any) is distributable to both operands.
632 //
633 // sext (add/sub nsw A, B) == add/sub nsw (sext A), (sext B)
634 // zext (add/sub nuw A, B) == add/sub nuw (zext A), (zext B)
635 if ((!SignExtended || BO->hasNoSignedWrap()) &&
636 (!ZeroExtended || BO->hasNoUnsignedWrap()))
637 return true;
638
639 if (BO->getOpcode() == Instruction::Add && !ZeroExtended && Idx) {
640 const auto *GEP = cast<GetElementPtrInst>(Val: Idx->getUser());
641 // For a sext(add nuw), allow tracing through when the enclosing GEP is both
642 // inbounds and nuw.
643 if (SignExtended && BO->hasNoUnsignedWrap() && GEP->isInBounds() &&
644 GEP->hasNoUnsignedWrap())
645 return true;
646
647 // Trace through sext when value tracking or the GEP's bounds prove that
648 // the addition cannot have signed overflow.
649 //
650 // Verified in @sext_add in split-gep.ll.
651 if (canReorderAddSextToGEP(Idx, Add: BO, DL))
652 return true;
653 }
654
655 return false;
656}
657
658std::optional<APInt> ConstantOffsetExtractor::findInOperands(RewriteNode &Node,
659 CastState &Casts) {
660 auto *BO = cast<BinaryOperator>(Val: Node.V);
661 // An intervening binary operator invalidates the GEP's index sign and bounds
662 // information, so do not pass it to either operand.
663 std::optional<APInt> LHS = find(V: BO->getOperand(i_nocapture: 0), Idx: nullptr, Casts);
664 if (LHS)
665 Node.LHS = UserTree.size() - 1;
666
667 std::optional<APInt> RHS = find(V: BO->getOperand(i_nocapture: 1), Idx: nullptr, Casts);
668 if (!RHS)
669 return LHS;
670 Node.RHS = UserTree.size() - 1;
671 if (BO->getOpcode() == Instruction::Sub)
672 *RHS = -*RHS;
673 return LHS ? *LHS + *RHS : *RHS;
674}
675
676std::optional<APInt> ConstantOffsetExtractor::find(Value *V, const Use *Idx,
677 CastState &Casts) {
678 // TODO(jingyue): We could trace into integer/pointer casts, such as
679 // inttoptr, ptrtoint, bitcast, and addrspacecast. We choose to handle only
680 // integers because it gives good enough results for our benchmarks.
681
682 // Bound repeated traversal of shared expressions and expansion of the
683 // rewrite tree. We can still extract offsets from the visited portion.
684 if (NumVisited++ >= MaxVisitedNodes)
685 return std::nullopt;
686
687 // We cannot do much with Values that are not a User, such as an Argument.
688 User *U = dyn_cast<User>(Val: V);
689 if (U == nullptr)
690 return std::nullopt;
691
692 RewriteNode Node{.V: U, .LHS: std::nullopt, .RHS: std::nullopt};
693 std::optional<APInt> ConstantOffset;
694 if (ConstantInt *CI = dyn_cast<ConstantInt>(Val: V)) {
695 // Leave literal zero offsets alone.
696 if (CI->isZero())
697 return std::nullopt;
698 // Hooray, we found it!
699 ConstantOffset = Casts.apply(Offset: CI->getValue());
700 Node.V = ConstantInt::getNullValue(Ty: CI->getType());
701 } else if (BinaryOperator *BO = dyn_cast<BinaryOperator>(Val: V)) {
702 // Trace into subexpressions for more hoisting opportunities.
703 if (canTraceInto(Casts, BO, Idx))
704 ConstantOffset = findInOperands(Node, Casts);
705 else if (BO->getOpcode() == Instruction::Xor) {
706 ConstantOffset = extractDisjointBitsFromXor(Node);
707 if (ConstantOffset)
708 *ConstantOffset = Casts.apply(Offset: *ConstantOffset);
709 }
710 } else if (isa<SExtInst, ZExtInst>(Val: V) ||
711 (isa<TruncInst>(Val: V) && !Casts.hasExtensions())) {
712 Casts.pushCast(Cast: cast<CastInst>(Val: V));
713 ConstantOffset = find(V: U->getOperand(i: 0), Idx, Casts);
714 Casts.popCast();
715 if (ConstantOffset)
716 Node.LHS = UserTree.size() - 1;
717 }
718
719 if (ConstantOffset)
720 UserTree.push_back(Elt: Node);
721 return ConstantOffset;
722}
723
724Value *ConstantOffsetExtractor::applyCasts(Value *V, const CastState &Casts) {
725 Value *Current = V;
726 // Casts is built in the use-def order. Therefore, we apply them to V in the
727 // reversed order.
728 for (CastInst *I : llvm::reverse(C: Casts.Casts)) {
729 if (Constant *C = dyn_cast<Constant>(Val: Current)) {
730 // Try to constant fold the cast.
731 Current = ConstantFoldCastOperand(Opcode: I->getOpcode(), C, DestTy: I->getType(), DL);
732 if (Current)
733 continue;
734 }
735
736 Instruction *Cast = I->clone();
737 Cast->setOperand(i: 0, Val: Current);
738 // In ConstantOffsetExtractor::find we do not analyze nuw/nsw for trunc, so
739 // we assume that it is ok to redistribute trunc over add/sub/or. But for
740 // example (add (trunc nuw A), (trunc nuw B)) would be more poisonous than
741 // (trunc nuw (add A, B))), and similarly, (sub (zext nneg A), (zext nneg
742 // B)) would be more poisonous than (zext nneg (sub nuw A, B)). To make such
743 // redistributions legal, we drop all the poison generating flags from the
744 // cloned cast instructions here.
745 Cast->dropPoisonGeneratingFlags();
746 Cast->insertBefore(BB&: *IP->getParent(), InsertPos: IP);
747 Current = Cast;
748 }
749 return Current;
750}
751
752Value *ConstantOffsetExtractor::rebuildWithoutConstOffset(unsigned NodeIndex,
753 CastState &Casts) {
754 const RewriteNode &Node = UserTree[NodeIndex];
755 if (isa<ConstantInt>(Val: Node.V))
756 return applyCasts(V: Node.V, Casts);
757
758 if (auto *Cast = dyn_cast<CastInst>(Val: Node.V)) {
759 Casts.pushCast(Cast);
760 Value *Result = rebuildWithoutConstOffset(NodeIndex: *Node.LHS, Casts);
761 Casts.popCast();
762 return Result;
763 }
764
765 auto *BO = cast<BinaryOperator>(Val: Node.V);
766 Value *LHS = Node.LHS ? rebuildWithoutConstOffset(NodeIndex: *Node.LHS, Casts)
767 : applyCasts(V: BO->getOperand(i_nocapture: 0), Casts);
768 Value *RHS = Node.RHS ? rebuildWithoutConstOffset(NodeIndex: *Node.RHS, Casts)
769 : applyCasts(V: BO->getOperand(i_nocapture: 1), Casts);
770
771 // Zero is a right identity for all supported operators, and a left identity
772 // for all except subtraction.
773 if (match(V: RHS, P: m_Zero()))
774 return LHS;
775 if (match(V: LHS, P: m_Zero()) && BO->getOpcode() != Instruction::Sub)
776 return RHS;
777
778 BinaryOperator::BinaryOps NewOp = BO->getOpcode();
779 if (BO->getOpcode() == Instruction::Or) {
780 // Rebuild "or" as "add", because "or" may be invalid for the new
781 // expression.
782 //
783 // For instance, given
784 // a | (b + 5) where a and b + 5 have no common bits,
785 // we can extract 5 as the constant offset.
786 //
787 // However, reusing the "or" in the new index would give us
788 // (a | b) + 5
789 // which does not equal a | (b + 5).
790 //
791 // Replacing the "or" with "add" is fine, because
792 // a | (b + 5) = a + (b + 5) = (a + b) + 5
793 NewOp = Instruction::Add;
794 }
795
796 BinaryOperator *NewBO = BinaryOperator::Create(Op: NewOp, S1: LHS, S2: RHS, Name: "", InsertBefore: IP);
797 NewBO->takeName(V: BO);
798 return NewBO;
799}
800
801std::optional<APInt>
802ConstantOffsetExtractor::extractDisjointBitsFromXor(RewriteNode &Node) {
803 auto *XorInst = cast<BinaryOperator>(Val: Node.V);
804 assert(XorInst && XorInst->getOpcode() == Instruction::Xor &&
805 "Expected XOR instruction");
806
807 Value *BaseOp;
808 ConstantInt *XorConstantOp;
809
810 if (!match(V: XorInst, P: m_Xor(L: m_Value(V&: BaseOp), R: m_ConstantInt(CI&: XorConstantOp))))
811 return std::nullopt;
812
813 const KnownBits BaseKnownBits = computeKnownBits(V: BaseOp, Q: SQ);
814 const APInt &ConstantValue = XorConstantOp->getValue();
815
816 // Compute the disjoint bits, i.e., those bits of the constant operand that
817 // are known-zero in the base. These disjoint bits will contribute to the
818 // final GEP offset. If there are no disjoint bits, there isn't any offset to
819 // extract from the xor.
820 const APInt DisjointBits = ConstantValue & BaseKnownBits.Zero;
821 if (DisjointBits.isZero())
822 return std::nullopt;
823
824 // Avoid a pessimizing rewrite if the disjoint bits include the sign bit.
825 if (DisjointBits.isSignBitSet())
826 return std::nullopt;
827
828 // Compute the remaining bits, i.e., the non-disjoint ones, which are those
829 // that must be preserved in the xor.
830 const APInt NonDisjointBits = ConstantValue & ~DisjointBits;
831 Node.RHS = UserTree.size();
832 UserTree.push_back(Elt: {.V: ConstantInt::get(Context&: XorInst->getContext(), V: NonDisjointBits),
833 .LHS: std::nullopt, .RHS: std::nullopt});
834
835 return DisjointBits;
836}
837
838/// Check if reassociating through a tree node would invalidate GEP's nuw flag.
839static bool allowsPreservingNUW(const User *U) {
840 if (const BinaryOperator *BO = dyn_cast<BinaryOperator>(Val: U)) {
841 // Binary operations need to be effectively add nuw.
842 auto Opcode = BO->getOpcode();
843 if (Opcode == BinaryOperator::Or) {
844 // Ors are only considered here if they are disjoint. The addition that
845 // they represent in this case is NUW.
846 assert(cast<PossiblyDisjointInst>(BO)->isDisjoint());
847 return true;
848 }
849 return Opcode == BinaryOperator::Add && BO->hasNoUnsignedWrap();
850 }
851 // UserTree can only contain ConstantInt, CastInst, or BinaryOperator.
852 // Among the possible CastInsts, only trunc without nuw is a problem: If it
853 // is distributed through an add nuw, wrapping may occur:
854 // "add nuw trunc(a), trunc(b)" is more poisonous than "trunc(add nuw a, b)"
855 if (const TruncInst *TI = dyn_cast<TruncInst>(Val: U))
856 return TI->hasNoUnsignedWrap();
857 assert((isa<CastInst>(U) || isa<ConstantInt>(U)) && "Unexpected User.");
858 return true;
859}
860
861static BasicBlock::iterator getIndexInsertionPoint(const Use &Idx) {
862 if (auto *I = dyn_cast<Instruction>(Val: Idx.get()))
863 if (auto IP = I->getInsertionPointAfterDef())
864 return *IP;
865 return cast<GetElementPtrInst>(Val: Idx.getUser())->getIterator();
866}
867
868Value *ConstantOffsetExtractor::Extract(const Use &Idx, User *&UserTreeRoot,
869 bool &PreservesNUW) {
870 ConstantOffsetExtractor Extractor(getIndexInsertionPoint(Idx));
871 // Find a constant offset first.
872 CastState Casts;
873 if (!Extractor.find(V: Idx, Idx: &Idx, Casts)) {
874 UserTreeRoot = nullptr;
875 PreservesNUW = true;
876 return nullptr;
877 }
878
879 PreservesNUW = all_of(Range&: Extractor.UserTree, P: [](const RewriteNode &Node) {
880 return allowsPreservingNUW(U: Node.V);
881 });
882
883 // Separates the constant offset from the GEP index.
884 Value *IdxWithoutConstOffset = Extractor.rebuildWithoutConstOffset();
885 UserTreeRoot = dyn_cast<User>(Val: IdxWithoutConstOffset);
886 return IdxWithoutConstOffset;
887}
888
889std::optional<APInt> ConstantOffsetExtractor::Find(const Use &Idx) {
890 auto *GEP = cast<GetElementPtrInst>(Val: Idx.getUser());
891 CastState Casts;
892 return ConstantOffsetExtractor(GEP->getIterator()).find(V: Idx, Idx: &Idx, Casts);
893}
894
895bool SeparateConstOffsetFromGEP::canonicalizeArrayIndicesToIndexSize(
896 GetElementPtrInst *GEP) {
897 bool Changed = false;
898 Type *PtrIdxTy = DL->getIndexType(PtrTy: GEP->getType());
899 gep_type_iterator GTI = gep_type_begin(GEP: *GEP);
900 for (User::op_iterator I = GEP->op_begin() + 1, E = GEP->op_end();
901 I != E; ++I, ++GTI) {
902 // Skip struct member indices which must be i32.
903 if (GTI.isSequential()) {
904 if ((*I)->getType() != PtrIdxTy) {
905 *I = CastInst::CreateIntegerCast(S: *I, Ty: PtrIdxTy, isSigned: true, Name: "idxprom",
906 InsertBefore: getIndexInsertionPoint(Idx: *I));
907 Changed = true;
908 }
909 }
910 }
911 return Changed;
912}
913
914APInt SeparateConstOffsetFromGEP::accumulateByteOffset(GetElementPtrInst *GEP,
915 bool &NeedsExtraction,
916 bool &SignedOverflow) {
917 NeedsExtraction = false;
918 SignedOverflow = false;
919 unsigned IdxWidth = DL->getIndexTypeSizeInBits(Ty: GEP->getType());
920 APInt AccumulativeByteOffset(IdxWidth, 0);
921 gep_type_iterator GTI = gep_type_begin(GEP: *GEP);
922 for (unsigned I = 1, E = GEP->getNumOperands(); I != E; ++I, ++GTI) {
923 if (GTI.isSequential()) {
924 // Constant offsets of scalable types are not really constant.
925 if (GTI.getIndexedType()->isScalableTy())
926 continue;
927
928 // Tries to extract a constant offset from this GEP index.
929 if (std::optional<APInt> ConstantOffset =
930 ConstantOffsetExtractor::Find(Idx: GEP->getOperandUse(i: I))) {
931 NeedsExtraction = true;
932 // A GEP may have multiple indices. We accumulate the extracted
933 // constant offset to a byte offset, and later offset the remainder of
934 // the original GEP with this byte offset.
935 const uint64_t Stride = GTI.getSequentialElementStride(DL: *DL);
936 bool Overflow;
937 auto ByteOffset = ConstantOffset->sextOrTrunc(width: IdxWidth).smul_ov(
938 RHS: APInt(IdxWidth, Stride, /*IsSigned=*/true, /*ImplicitTrunc=*/true),
939 Overflow);
940 SignedOverflow |= Overflow;
941 // Removing an offset must not introduce signed overflow. For strides
942 // greater than one, the non-wrapping multiplication of both the
943 // original index and extracted offset already rules out this wrap.
944 if (Stride == 1) {
945 auto *Offset = ConstantInt::get(Context&: GEP->getContext(), V: *ConstantOffset);
946 SignedOverflow |=
947 computeOverflowForSignedSub(LHS: GEP->getOperand(i_nocapture: I), RHS: Offset, SQ: *DL) !=
948 OverflowResult::NeverOverflows;
949 }
950 AccumulativeByteOffset =
951 AccumulativeByteOffset.sadd_ov(RHS: ByteOffset, Overflow);
952 SignedOverflow |= Overflow;
953 }
954 } else if (LowerGEP) {
955 StructType *StTy = GTI.getStructType();
956 uint64_t Field = cast<ConstantInt>(Val: GEP->getOperand(i_nocapture: I))->getZExtValue();
957 // Skip field 0 as the offset is always 0.
958 if (Field != 0) {
959 NeedsExtraction = true;
960 AccumulativeByteOffset +=
961 APInt(IdxWidth, DL->getStructLayout(Ty: StTy)->getElementOffset(Idx: Field),
962 /*IsSigned=*/true, /*ImplicitTrunc=*/true);
963 }
964 }
965 }
966 return AccumulativeByteOffset;
967}
968
969void SeparateConstOffsetFromGEP::lowerToSingleIndexGEPs(
970 GetElementPtrInst *Variadic, const APInt &AccumulativeByteOffset) {
971 IRBuilder<> Builder(Variadic);
972 Type *PtrIndexTy = DL->getIndexType(PtrTy: Variadic->getType());
973
974 Value *ResultPtr = Variadic->getOperand(i_nocapture: 0);
975 Loop *L = LI->getLoopFor(BB: Variadic->getParent());
976 // Check if the base is not loop invariant or used more than once.
977 bool isSwapCandidate =
978 L && L->isLoopInvariant(V: ResultPtr) &&
979 !hasMoreThanOneUseInLoop(v: ResultPtr, L);
980 Value *FirstResult = nullptr;
981
982 gep_type_iterator GTI = gep_type_begin(GEP: *Variadic);
983 // Create an ugly GEP for each sequential index. We don't create GEPs for
984 // structure indices, as they are accumulated in the constant offset index.
985 for (unsigned I = 1, E = Variadic->getNumOperands(); I != E; ++I, ++GTI) {
986 if (GTI.isSequential()) {
987 Value *Idx = Variadic->getOperand(i_nocapture: I);
988 // Skip zero indices.
989 if (ConstantInt *CI = dyn_cast<ConstantInt>(Val: Idx))
990 if (CI->isZero())
991 continue;
992
993 APInt ElementSize = APInt(PtrIndexTy->getIntegerBitWidth(),
994 GTI.getSequentialElementStride(DL: *DL));
995 // Scale the index by element size.
996 if (ElementSize != 1) {
997 if (ElementSize.isPowerOf2()) {
998 Idx = Builder.CreateShl(
999 LHS: Idx, RHS: ConstantInt::get(Ty: PtrIndexTy, V: ElementSize.logBase2()));
1000 } else {
1001 Idx =
1002 Builder.CreateMul(LHS: Idx, RHS: ConstantInt::get(Ty: PtrIndexTy, V: ElementSize));
1003 }
1004 }
1005 // Create an ugly GEP with a single index for each index.
1006 ResultPtr = Builder.CreatePtrAdd(Ptr: ResultPtr, Offset: Idx, Name: "uglygep");
1007 if (FirstResult == nullptr)
1008 FirstResult = ResultPtr;
1009 }
1010 }
1011
1012 // Create a GEP with the constant offset index.
1013 if (AccumulativeByteOffset != 0) {
1014 Value *Offset = ConstantInt::get(Ty: PtrIndexTy, V: AccumulativeByteOffset);
1015 ResultPtr = Builder.CreatePtrAdd(Ptr: ResultPtr, Offset, Name: "uglygep");
1016 } else
1017 isSwapCandidate = false;
1018
1019 // If we created a GEP with constant index, and the base is loop invariant,
1020 // then we swap the first one with it, so LICM can move constant GEP out
1021 // later.
1022 auto *FirstGEP = dyn_cast_or_null<GetElementPtrInst>(Val: FirstResult);
1023 auto *SecondGEP = dyn_cast<GetElementPtrInst>(Val: ResultPtr);
1024 if (isSwapCandidate && isLegalToSwapOperand(First: FirstGEP, Second: SecondGEP, CurLoop: L))
1025 swapGEPOperand(First: FirstGEP, Second: SecondGEP);
1026
1027 Variadic->replaceAllUsesWith(V: ResultPtr);
1028 Variadic->eraseFromParent();
1029}
1030
1031bool SeparateConstOffsetFromGEP::reorderGEP(GetElementPtrInst *GEP,
1032 TargetTransformInfo &TTI) {
1033 auto PtrGEP = dyn_cast<GetElementPtrInst>(Val: GEP->getPointerOperand());
1034 if (!PtrGEP)
1035 return false;
1036
1037 bool NestedNeedsExtraction, OffsetOverflow;
1038 APInt NestedByteOffset =
1039 accumulateByteOffset(GEP: PtrGEP, NeedsExtraction&: NestedNeedsExtraction, SignedOverflow&: OffsetOverflow);
1040 if (!NestedNeedsExtraction)
1041 return false;
1042
1043 unsigned AddrSpace = PtrGEP->getPointerAddressSpace();
1044 if (!TTI.isLegalAddressingMode(Ty: GEP->getResultElementType(),
1045 /*BaseGV=*/nullptr,
1046 BaseOffset: NestedByteOffset.getSExtValue(),
1047 /*HasBaseReg=*/true, /*Scale=*/0, AddrSpace))
1048 return false;
1049
1050 bool GEPInBounds = GEP->isInBounds();
1051 bool PtrGEPInBounds = PtrGEP->isInBounds();
1052 bool IsChainInBounds = GEPInBounds && PtrGEPInBounds;
1053 if (IsChainInBounds) {
1054 auto IsKnownNonNegative = [this](Value *V) {
1055 return isKnownNonNegative(V, SQ: *DL);
1056 };
1057 IsChainInBounds &= all_of(Range: GEP->indices(), P: IsKnownNonNegative);
1058 if (IsChainInBounds)
1059 IsChainInBounds &= all_of(Range: PtrGEP->indices(), P: IsKnownNonNegative);
1060 }
1061
1062 IRBuilder<> Builder(GEP);
1063 // For trivial GEP chains, we can swap the indices.
1064 Value *NewSrc = Builder.CreateGEP(
1065 Ty: GEP->getSourceElementType(), Ptr: PtrGEP->getPointerOperand(),
1066 IdxList: SmallVector<Value *, 4>(GEP->indices()), Name: "", NW: IsChainInBounds);
1067 Value *NewGEP = Builder.CreateGEP(Ty: PtrGEP->getSourceElementType(), Ptr: NewSrc,
1068 IdxList: SmallVector<Value *, 4>(PtrGEP->indices()),
1069 Name: "", NW: IsChainInBounds);
1070 GEP->replaceAllUsesWith(V: NewGEP);
1071 RecursivelyDeleteTriviallyDeadInstructions(V: GEP);
1072 return true;
1073}
1074
1075bool SeparateConstOffsetFromGEP::splitGEP(GetElementPtrInst *GEP) {
1076 // Skip vector GEPs.
1077 if (GEP->getType()->isVectorTy())
1078 return false;
1079
1080 // If the base of this GEP is a ptradd of a constant, lets pass the constant
1081 // along. This ensures that when we have a chain of GEPs the constant
1082 // offset from each is accumulated.
1083 Value *NewBase;
1084 const APInt *BaseOffset;
1085 bool ExtractBase = match(V: GEP->getPointerOperand(),
1086 P: m_PtrAdd(PointerOp: m_Value(V&: NewBase), OffsetOp: m_APInt(Res&: BaseOffset)));
1087
1088 unsigned IdxWidth = DL->getIndexTypeSizeInBits(Ty: GEP->getType());
1089 APInt BaseByteOffset =
1090 ExtractBase ? BaseOffset->sextOrTrunc(width: IdxWidth) : APInt(IdxWidth, 0);
1091
1092 // The backend can already nicely handle the case where all indices are
1093 // constant.
1094 if (GEP->hasAllConstantIndices() && !ExtractBase)
1095 return false;
1096
1097 bool Changed = canonicalizeArrayIndicesToIndexSize(GEP);
1098
1099 bool NeedsExtraction, OffsetOverflow;
1100 APInt NonBaseByteOffset =
1101 accumulateByteOffset(GEP, NeedsExtraction, SignedOverflow&: OffsetOverflow);
1102 bool AddOverflow;
1103 APInt AccumulativeByteOffset =
1104 BaseByteOffset.sadd_ov(RHS: NonBaseByteOffset, Overflow&: AddOverflow);
1105 OffsetOverflow |= AddOverflow;
1106
1107 TargetTransformInfo &TTI = GetTTI(*GEP->getFunction());
1108
1109 if (!NeedsExtraction && !ExtractBase) {
1110 Changed |= reorderGEP(GEP, TTI);
1111 return Changed;
1112 }
1113
1114 // If LowerGEP is disabled, before really splitting the GEP, check whether the
1115 // backend supports the addressing mode we are about to produce. If no, this
1116 // splitting probably won't be beneficial.
1117 // If LowerGEP is enabled, even the extracted constant offset can not match
1118 // the addressing mode, we can still do optimizations to other lowered parts
1119 // of variable indices. Therefore, we don't check for addressing modes in that
1120 // case.
1121 if (!LowerGEP) {
1122 unsigned AddrSpace = GEP->getPointerAddressSpace();
1123 if (!TTI.isLegalAddressingMode(
1124 Ty: GEP->getResultElementType(),
1125 /*BaseGV=*/nullptr, BaseOffset: AccumulativeByteOffset.getSExtValue(),
1126 /*HasBaseReg=*/true, /*Scale=*/0, AddrSpace)) {
1127 // If the addressing mode was not legal and the base byte offset was not
1128 // 0, it could be a case where the total offset became too large for
1129 // the addressing mode. Try again without extracting the base offset.
1130 if (!ExtractBase)
1131 return Changed;
1132 ExtractBase = false;
1133 BaseByteOffset = APInt(IdxWidth, 0);
1134 AccumulativeByteOffset = NonBaseByteOffset;
1135 if (!TTI.isLegalAddressingMode(
1136 Ty: GEP->getResultElementType(),
1137 /*BaseGV=*/nullptr, BaseOffset: AccumulativeByteOffset.getSExtValue(),
1138 /*HasBaseReg=*/true, /*Scale=*/0, AddrSpace))
1139 return Changed;
1140 // We can proceed with just extracting the other (non-base) offsets.
1141 NeedsExtraction = true;
1142 }
1143 }
1144
1145 // Track information for preserving GEP flags.
1146 bool AllOffsetsNonNegative =
1147 AccumulativeByteOffset.isNonNegative() && !OffsetOverflow;
1148 bool AllNUWPreserved = GEP->hasNoUnsignedWrap();
1149 bool NewGEPInBounds = GEP->isInBounds();
1150 bool NewGEPNUSW = GEP->hasNoUnsignedSignedWrap();
1151
1152 // Remove the constant offset in each sequential index. The resultant GEP
1153 // computes the variadic base.
1154 // Notice that we don't remove struct field indices here. If LowerGEP is
1155 // disabled, a structure index is not accumulated and we still use the old
1156 // one. If LowerGEP is enabled, a structure index is accumulated in the
1157 // constant offset. LowerToSingleIndexGEPs will later handle the constant
1158 // offset and won't need a new structure index.
1159 gep_type_iterator GTI = gep_type_begin(GEP: *GEP);
1160 for (unsigned I = 1, E = GEP->getNumOperands(); I != E; ++I, ++GTI) {
1161 if (GTI.isSequential()) {
1162 // Constant offsets of scalable types are not really constant.
1163 if (GTI.getIndexedType()->isScalableTy())
1164 continue;
1165
1166 // Splits this GEP index into a variadic part and a constant offset, and
1167 // uses the variadic part as the new index.
1168 Value *Idx = GEP->getOperand(i_nocapture: I);
1169 User *UserTreeRoot;
1170 bool PreservesNUW;
1171 Value *NewIdx = ConstantOffsetExtractor::Extract(
1172 Idx: GEP->getOperandUse(i: I), UserTreeRoot, PreservesNUW);
1173 if (NewIdx != nullptr) {
1174 // Switches to the index with the constant offset removed.
1175 GEP->setOperand(i_nocapture: I, Val_nocapture: NewIdx);
1176 // After switching to the new index, we can garbage-collect the
1177 // rewritten tree and the old index if they are not used.
1178 if (auto *I = dyn_cast_or_null<Instruction>(Val: UserTreeRoot))
1179 RecursivelyDeleteTriviallyDeadInstructions(V: I);
1180 if (auto *I = dyn_cast<Instruction>(Val: Idx))
1181 RecursivelyDeleteTriviallyDeadInstructions(V: I);
1182 Idx = NewIdx;
1183 AllNUWPreserved &= PreservesNUW;
1184 }
1185 AllOffsetsNonNegative =
1186 AllOffsetsNonNegative && isKnownNonNegative(V: Idx, SQ: *DL);
1187 }
1188 }
1189 if (ExtractBase) {
1190 GEPOperator *Base = cast<GEPOperator>(Val: GEP->getPointerOperand());
1191 AllNUWPreserved &= Base->hasNoUnsignedWrap();
1192 NewGEPInBounds &= Base->isInBounds();
1193 NewGEPNUSW &= Base->hasNoUnsignedSignedWrap();
1194 AllOffsetsNonNegative &= BaseByteOffset.isNonNegative();
1195
1196 GEP->setOperand(i_nocapture: 0, Val_nocapture: NewBase);
1197 RecursivelyDeleteTriviallyDeadInstructions(V: Base);
1198 }
1199
1200 // Clear the inbounds attribute because the new index may be off-bound.
1201 // e.g.,
1202 //
1203 // b = add i64 a, 5
1204 // addr = gep inbounds float, float* p, i64 b
1205 //
1206 // is transformed to:
1207 //
1208 // addr2 = gep float, float* p, i64 a ; inbounds removed
1209 // addr = gep float, float* addr2, i64 5 ; inbounds removed
1210 //
1211 // If a is -4, although the old index b is in bounds, the new index a is
1212 // off-bound. http://llvm.org/docs/LangRef.html#id181 says "if the
1213 // inbounds keyword is not present, the offsets are added to the base
1214 // address with silently-wrapping two's complement arithmetic".
1215 // Therefore, the final code will be a semantically equivalent.
1216 GEPNoWrapFlags NewGEPFlags = GEPNoWrapFlags::none();
1217
1218 // If removing and scaling offsets cannot have signed overflow, and all
1219 // remaining indices and accumulated offsets are nonnegative, they can be
1220 // added in any order. An initially inbounds/nusw GEP then guarantees that the
1221 // intermediate results stay in bounds and do not overflow in a nusw sense.
1222 bool CanPreserveInBoundsNUSW = AllOffsetsNonNegative;
1223
1224 // If the initial GEP was NUW and all operations that we reassociate were NUW
1225 // additions, the resulting GEPs are also NUW.
1226 if (AllNUWPreserved) {
1227 NewGEPFlags |= GEPNoWrapFlags::noUnsignedWrap();
1228 // If the initial GEP additionally had NUSW (or inbounds, which implies
1229 // NUSW), we know that the indices in the initial GEP must all have their
1230 // signbit not set. For indices that are the result of NUW adds, the
1231 // add-operands therefore also don't have their signbit set. Therefore, all
1232 // indices of the resulting GEPs are non-negative -> we can preserve
1233 // the inbounds/nusw flag.
1234 CanPreserveInBoundsNUSW |= NewGEPNUSW;
1235 }
1236
1237 if (CanPreserveInBoundsNUSW) {
1238 if (NewGEPInBounds)
1239 NewGEPFlags |= GEPNoWrapFlags::inBounds();
1240 else if (NewGEPNUSW)
1241 NewGEPFlags |= GEPNoWrapFlags::noUnsignedSignedWrap();
1242 }
1243
1244 GEP->setNoWrapFlags(NewGEPFlags);
1245
1246 // Lowers a GEP to GEPs with a single index.
1247 if (LowerGEP) {
1248 lowerToSingleIndexGEPs(Variadic: GEP, AccumulativeByteOffset);
1249 return true;
1250 }
1251
1252 // No need to create another GEP if the accumulative byte offset is 0.
1253 if (AccumulativeByteOffset == 0)
1254 return true;
1255
1256 // Offsets the base with the accumulative byte offset.
1257 //
1258 // %gep ; the base
1259 // ... %gep ...
1260 //
1261 // => add the offset
1262 //
1263 // %gep2 ; clone of %gep
1264 // %new.gep = gep i8, %gep2, %offset
1265 // %gep ; will be removed
1266 // ... %gep ...
1267 //
1268 // => replace all uses of %gep with %new.gep and remove %gep
1269 //
1270 // %gep2 ; clone of %gep
1271 // %new.gep = gep i8, %gep2, %offset
1272 // ... %new.gep ...
1273 Instruction *NewGEP = GEP->clone();
1274 NewGEP->insertBefore(InsertPos: GEP->getIterator());
1275
1276 Type *PtrIdxTy = DL->getIndexType(PtrTy: GEP->getType());
1277 IRBuilder<> Builder(GEP);
1278 NewGEP = cast<Instruction>(Val: Builder.CreatePtrAdd(
1279 Ptr: NewGEP, Offset: ConstantInt::get(Ty: PtrIdxTy, V: AccumulativeByteOffset),
1280 Name: GEP->getName(), NW: NewGEPFlags));
1281 NewGEP->copyMetadata(SrcInst: *GEP);
1282
1283 GEP->replaceAllUsesWith(V: NewGEP);
1284 GEP->eraseFromParent();
1285
1286 return true;
1287}
1288
1289bool SeparateConstOffsetFromGEPLegacyPass::runOnFunction(Function &F) {
1290 if (skipFunction(F))
1291 return false;
1292 auto *DT = &getAnalysis<DominatorTreeWrapperPass>().getDomTree();
1293 auto *LI = &getAnalysis<LoopInfoWrapperPass>().getLoopInfo();
1294 auto *TLI = &getAnalysis<TargetLibraryInfoWrapperPass>().getTLI(F);
1295 auto GetTTI = [this](Function &F) -> TargetTransformInfo & {
1296 return this->getAnalysis<TargetTransformInfoWrapperPass>().getTTI(F);
1297 };
1298 SeparateConstOffsetFromGEP Impl(DT, LI, TLI, GetTTI, LowerGEP);
1299 return Impl.run(F);
1300}
1301
1302bool SeparateConstOffsetFromGEP::run(Function &F) {
1303 const ScalarOptions &Opts =
1304 F.getContext().getLibraryOptions().get<ScalarOptions>();
1305 if (Opts.disable_separate_const_offset_from_gep)
1306 return false;
1307
1308 DL = &F.getDataLayout();
1309 bool Changed = false;
1310
1311 ReversePostOrderTraversal<Function *> RPOT(&F);
1312 for (BasicBlock *B : RPOT) {
1313 if (!DT->isReachableFromEntry(A: B))
1314 continue;
1315
1316 for (Instruction &I : llvm::make_early_inc_range(Range&: *B))
1317 if (GetElementPtrInst *GEP = dyn_cast<GetElementPtrInst>(Val: &I))
1318 Changed |= splitGEP(GEP);
1319 // No need to split GEP ConstantExprs because all its indices are constant
1320 // already.
1321 }
1322
1323 Changed |= reuniteExts(F);
1324
1325 if (Opts.reassociate_geps_verify_no_dead_code)
1326 verifyNoDeadCode(F);
1327
1328 return Changed;
1329}
1330
1331Instruction *SeparateConstOffsetFromGEP::findClosestMatchingDominator(
1332 ExprKey Key, Instruction *Dominatee,
1333 DenseMap<ExprKey, SmallVector<Instruction *, 2>> &DominatingExprs) {
1334 auto Pos = DominatingExprs.find(Val: Key);
1335 if (Pos == DominatingExprs.end())
1336 return nullptr;
1337
1338 auto &Candidates = Pos->second;
1339 // Because we process the basic blocks in pre-order of the dominator tree, a
1340 // candidate that doesn't dominate the current instruction won't dominate any
1341 // future instruction either. Therefore, we pop it out of the stack. This
1342 // optimization makes the algorithm O(n).
1343 while (!Candidates.empty()) {
1344 Instruction *Candidate = Candidates.back();
1345 if (DT->dominates(Def: Candidate, User: Dominatee))
1346 return Candidate;
1347 Candidates.pop_back();
1348 }
1349 return nullptr;
1350}
1351
1352bool SeparateConstOffsetFromGEP::reuniteExts(Instruction *I) {
1353 if (!I->getType()->isIntOrIntVectorTy())
1354 return false;
1355
1356 // Dom: LHS+RHS
1357 // I: sext(LHS)+sext(RHS)
1358 // If Dom can't sign overflow and Dom dominates I, optimize I to sext(Dom).
1359 // TODO: handle zext
1360 Value *LHS = nullptr, *RHS = nullptr;
1361 if (match(V: I, P: m_Add(L: m_SExt(Op: m_Value(V&: LHS)), R: m_SExt(Op: m_Value(V&: RHS))))) {
1362 if (LHS->getType() == RHS->getType()) {
1363 ExprKey Key = createNormalizedCommutablePair(A: LHS, B: RHS);
1364 if (auto *Dom = findClosestMatchingDominator(Key, Dominatee: I, DominatingExprs&: DominatingAdds)) {
1365 Instruction *NewSExt =
1366 new SExtInst(Dom, I->getType(), "", I->getIterator());
1367 NewSExt->takeName(V: I);
1368 I->replaceAllUsesWith(V: NewSExt);
1369 NewSExt->setDebugLoc(I->getDebugLoc());
1370 RecursivelyDeleteTriviallyDeadInstructions(V: I);
1371 return true;
1372 }
1373 }
1374 } else if (match(V: I, P: m_Sub(L: m_SExt(Op: m_Value(V&: LHS)), R: m_SExt(Op: m_Value(V&: RHS))))) {
1375 if (LHS->getType() == RHS->getType()) {
1376 if (auto *Dom =
1377 findClosestMatchingDominator(Key: {LHS, RHS}, Dominatee: I, DominatingExprs&: DominatingSubs)) {
1378 Instruction *NewSExt =
1379 new SExtInst(Dom, I->getType(), "", I->getIterator());
1380 NewSExt->takeName(V: I);
1381 I->replaceAllUsesWith(V: NewSExt);
1382 NewSExt->setDebugLoc(I->getDebugLoc());
1383 RecursivelyDeleteTriviallyDeadInstructions(V: I);
1384 return true;
1385 }
1386 }
1387 }
1388
1389 // Add I to DominatingExprs if it's an add/sub that can't sign overflow.
1390 if (match(V: I, P: m_NSWAdd(L: m_Value(V&: LHS), R: m_Value(V&: RHS)))) {
1391 if (programUndefinedIfPoison(Inst: I)) {
1392 ExprKey Key = createNormalizedCommutablePair(A: LHS, B: RHS);
1393 DominatingAdds[Key].push_back(Elt: I);
1394 }
1395 } else if (match(V: I, P: m_NSWSub(L: m_Value(V&: LHS), R: m_Value(V&: RHS)))) {
1396 if (programUndefinedIfPoison(Inst: I))
1397 DominatingSubs[{LHS, RHS}].push_back(Elt: I);
1398 }
1399 return false;
1400}
1401
1402bool SeparateConstOffsetFromGEP::reuniteExts(Function &F) {
1403 bool Changed = false;
1404 DominatingAdds.clear();
1405 DominatingSubs.clear();
1406 for (const auto Node : depth_first(G: DT)) {
1407 BasicBlock *BB = Node->getBlock();
1408 for (Instruction &I : llvm::make_early_inc_range(Range&: *BB))
1409 Changed |= reuniteExts(I: &I);
1410 }
1411 return Changed;
1412}
1413
1414void SeparateConstOffsetFromGEP::verifyNoDeadCode(Function &F) {
1415 for (BasicBlock &B : F) {
1416 for (Instruction &I : B) {
1417 if (isInstructionTriviallyDead(I: &I)) {
1418 std::string ErrMessage;
1419 raw_string_ostream RSO(ErrMessage);
1420 RSO << "Dead instruction detected!\n" << I << "\n";
1421 llvm_unreachable(RSO.str().c_str());
1422 }
1423 }
1424 }
1425}
1426
1427bool SeparateConstOffsetFromGEP::isLegalToSwapOperand(
1428 GetElementPtrInst *FirstGEP, GetElementPtrInst *SecondGEP, Loop *CurLoop) {
1429 if (!FirstGEP || !FirstGEP->hasOneUse())
1430 return false;
1431
1432 if (!SecondGEP || FirstGEP->getParent() != SecondGEP->getParent())
1433 return false;
1434
1435 if (FirstGEP == SecondGEP)
1436 return false;
1437
1438 unsigned FirstNum = FirstGEP->getNumOperands();
1439 unsigned SecondNum = SecondGEP->getNumOperands();
1440 // Give up if the number of operands are not 2.
1441 if (FirstNum != SecondNum || FirstNum != 2)
1442 return false;
1443
1444 Value *FirstBase = FirstGEP->getOperand(i_nocapture: 0);
1445 Value *SecondBase = SecondGEP->getOperand(i_nocapture: 0);
1446 Value *FirstOffset = FirstGEP->getOperand(i_nocapture: 1);
1447 // Give up if the index of the first GEP is loop invariant.
1448 if (CurLoop->isLoopInvariant(V: FirstOffset))
1449 return false;
1450
1451 // Give up if base doesn't have same type.
1452 if (FirstBase->getType() != SecondBase->getType())
1453 return false;
1454
1455 Instruction *FirstOffsetDef = dyn_cast<Instruction>(Val: FirstOffset);
1456
1457 // Check if the second operand of first GEP has constant coefficient.
1458 // For an example, for the following code, we won't gain anything by
1459 // hoisting the second GEP out because the second GEP can be folded away.
1460 // %scevgep.sum.ur159 = add i64 %idxprom48.ur, 256
1461 // %67 = shl i64 %scevgep.sum.ur159, 2
1462 // %uglygep160 = getelementptr i8* %65, i64 %67
1463 // %uglygep161 = getelementptr i8* %uglygep160, i64 -1024
1464
1465 // Skip constant shift instruction which may be generated by Splitting GEPs.
1466 if (FirstOffsetDef && FirstOffsetDef->isShift() &&
1467 isa<ConstantInt>(Val: FirstOffsetDef->getOperand(i: 1)))
1468 FirstOffsetDef = dyn_cast<Instruction>(Val: FirstOffsetDef->getOperand(i: 0));
1469
1470 // Give up if FirstOffsetDef is an Add or Sub with constant.
1471 // Because it may not profitable at all due to constant folding.
1472 if (FirstOffsetDef)
1473 if (BinaryOperator *BO = dyn_cast<BinaryOperator>(Val: FirstOffsetDef)) {
1474 unsigned opc = BO->getOpcode();
1475 if ((opc == Instruction::Add || opc == Instruction::Sub) &&
1476 (isa<ConstantInt>(Val: BO->getOperand(i_nocapture: 0)) ||
1477 isa<ConstantInt>(Val: BO->getOperand(i_nocapture: 1))))
1478 return false;
1479 }
1480 return true;
1481}
1482
1483bool SeparateConstOffsetFromGEP::hasMoreThanOneUseInLoop(Value *V, Loop *L) {
1484 // TODO: Could look at uses of globals, but we need to make sure we are
1485 // looking at the correct function.
1486 if (isa<Constant>(Val: V))
1487 return false;
1488
1489 int UsesInLoop = 0;
1490 for (User *U : V->users()) {
1491 if (Instruction *User = dyn_cast<Instruction>(Val: U))
1492 if (L->contains(Inst: User))
1493 if (++UsesInLoop > 1)
1494 return true;
1495 }
1496 return false;
1497}
1498
1499void SeparateConstOffsetFromGEP::swapGEPOperand(GetElementPtrInst *First,
1500 GetElementPtrInst *Second) {
1501 Value *Offset1 = First->getOperand(i_nocapture: 1);
1502 Value *Offset2 = Second->getOperand(i_nocapture: 1);
1503 First->setOperand(i_nocapture: 1, Val_nocapture: Offset2);
1504 Second->setOperand(i_nocapture: 1, Val_nocapture: Offset1);
1505
1506 // After changing (p+o)+c to (p+c)+o, the inner GEP may not be inbounds
1507 // anymore.
1508 const DataLayout &DAL = First->getDataLayout();
1509 unsigned IdxBits = DAL.getIndexSizeInBits(
1510 AS: cast<PointerType>(Val: First->getType())->getAddressSpace());
1511
1512 auto ClearNoWrapFlags = [&] {
1513 // TODO(gep_nowrap): Make flag preservation more precise.
1514 First->setNoWrapFlags(GEPNoWrapFlags::none());
1515 Second->setNoWrapFlags(GEPNoWrapFlags::none());
1516 };
1517
1518 APInt FirstOffset(IdxBits, 0);
1519 if (!First->accumulateConstantOffset(DL: DAL, Offset&: FirstOffset)) {
1520 ClearNoWrapFlags();
1521 return;
1522 }
1523
1524 APInt BaseOffset(IdxBits, 0);
1525 Value *NewBase =
1526 First->getOperand(i_nocapture: 0)->stripAndAccumulateInBoundsConstantOffsets(
1527 DL: DAL, Offset&: BaseOffset);
1528
1529 bool Overflow = false;
1530 APInt TotalOffset = BaseOffset.uadd_ov(RHS: FirstOffset, Overflow);
1531 uint64_t ObjectSize;
1532 if (Overflow || !getObjectSize(Ptr: NewBase, Size&: ObjectSize, DL: DAL, TLI) ||
1533 TotalOffset.ugt(RHS: ObjectSize)) {
1534 ClearNoWrapFlags();
1535 return;
1536 }
1537
1538 First->setIsInBounds(true);
1539}
1540
1541void SeparateConstOffsetFromGEPPass::printPipeline(
1542 raw_ostream &OS, function_ref<StringRef(StringRef)> MapClassName2PassName) {
1543 static_cast<PassInfoMixin<SeparateConstOffsetFromGEPPass> *>(this)
1544 ->printPipeline(OS, MapClassName2PassName);
1545 OS << '<';
1546 if (LowerGEP)
1547 OS << "lower-gep";
1548 OS << '>';
1549}
1550
1551PreservedAnalyses
1552SeparateConstOffsetFromGEPPass::run(Function &F, FunctionAnalysisManager &AM) {
1553 auto *DT = &AM.getResult<DominatorTreeAnalysis>(IR&: F);
1554 auto *LI = &AM.getResult<LoopAnalysis>(IR&: F);
1555 auto *TLI = &AM.getResult<TargetLibraryAnalysis>(IR&: F);
1556 auto GetTTI = [&AM](Function &F) -> TargetTransformInfo & {
1557 return AM.getResult<TargetIRAnalysis>(IR&: F);
1558 };
1559 SeparateConstOffsetFromGEP Impl(DT, LI, TLI, GetTTI, LowerGEP);
1560 if (!Impl.run(F))
1561 return PreservedAnalyses::all();
1562 PreservedAnalyses PA;
1563 PA.preserveSet<CFGAnalyses>();
1564 return PA;
1565}
1566