[SeparateConstOffsetFromGEP] inbounds zext => sext for better splitting

[jingyue]

Summary

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

Depends on D4056

[meheff]

LGTM after clarifying and fixing some of the comments.

lib/Transforms/Scalar/SeparateConstOffsetFromGEP.cpp

268

spelling: Verified. Also it would be good to include in this comment why it is advantageous to convert zexts in to sexts. It's not immediately clear because like zext for some values of a and b: sext(a + b) != sext(a) + sext(b). Include the how sext and add with positive result is handled as a special case while no such case exists for zext.

316

Doesn't have to be done in this change, but this function would be easier to follow if every if statement conditionally returned true and then false was returned at the bottom (or vice versa). That is, the if statements should explicitly enumerate the accepted cases. As it is, some return true some return false which implies some dependencies between the ifs which makes the logic more complicated than necessary to follow.

621

This comment is a little confusing. Maybe something like:

For GEP operand a, if a <= max signed value of typeof(a), then the sign bit of a is zero and sext(a) = zext(a). Because the GEP is inbounds, we know a <= ObjectSize, so the necessary condition can be reduced to ObjectSize <= max signed value of typeof(a).

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

102

Similar to sum_of_array2

[jingyue]

Thanks! Updated the comments.

[jingyue]

lib/Transforms/Scalar/SeparateConstOffsetFromGEP.cpp

268

done

316

ack'ed

621

done

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

102

actually more similar to sum_of_array

[jingyue]

Committed in r210444

lib/Transforms/Scalar/SeparateConstOffsetFromGEP.cpp

public:
   SeparateConstOffsetFromGEP() : FunctionPass(ID) {
     initializeSeparateConstOffsetFromGEPPass(*PassRegistry::getPassRegistry());
   }
 
   void getAnalysisUsage(AnalysisUsage &AU) const override {
     AU.addRequired<DataLayoutPass>();
     AU.addRequired<TargetTransformInfo>();
   }
+
+  bool doInitialization(Module &M) override {
+    DataLayoutPass *DLP = getAnalysisIfAvailable<DataLayoutPass>();
+    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, const DataLayout *DL,
-                               bool &NeedsExtraction);
+  int64_t accumulateByteOffset(GetElementPtrInst *GEP, bool &NeedsExtraction);
+  /// Canonicalize array indices to pointer-size integers. This helps to
+  /// simplify the logic of splitting a GEP. For example, if a + b is a
+  /// pointer-size integer, we have
+  ///   gep base, a + b = gep (gep base, a), b
+  /// However, this equality may not hold if the size of a + b is smaller than
+  /// the pointer size, because LLVM conceptually sign-extends GEP indices to
+  /// pointer size before computing the address
+  /// (http://llvm.org/docs/LangRef.html#id181).
+  ///
+  /// This canonicalization is very likely already done in clang and
+  /// instcombine. Therefore, the program will probably remain the same.
+  ///
+  /// Returns true if the module changes.
+  ///
+  /// Verified in @i32_add in split-gep.ll
+  bool canonicalizeArrayIndicesToPointerSize(GetElementPtrInst *GEP);
+  /// For each array index that is in the form of zext(a), convert it to sext(a)
+  /// if we can prove zext(a) <= max signed value of typeof(a). We prefer
+  /// sext(a) to zext(a), because in the special case where x + y >= 0 and
+  /// (x >= 0 or y >= 0), function CanTraceInto can split sext(x + y),
+  /// while no such case exists for zext(x + y).
+  ///
+  /// Note that
+  ///   zext(x + y) = zext(x) + zext(y)
+  /// is wrong, e.g.,
+  ///   zext i32(UINT_MAX + 1) to i64 !=
+  ///   (zext i32 UINT_MAX to i64) + (zext i32 1 to i64)
+  ///
+  /// Returns true if the module changes.
+  ///
+  /// Verified in @inbounds_zext_add in split-gep.ll and @sum_of_array3 in
+  /// split-gep-and-gvn.ll
+  bool convertInBoundsZExtToSExt(GetElementPtrInst *GEP);
+
+  const DataLayout *DL;
 };
 }  // anonymous namespace
 
 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;
   }
bool ConstantOffsetExtractor::CanTraceInto(bool SignExtended,
   // 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 && NonNegative) {
+  if (BO->getOpcode() == Instruction::Add && !ZeroExtended && NonNegative) {
     // If a + b >= 0 and (a >= 0 or b >= 0), then
-    //   s/zext(a + b) = s/zext(a) + s/zext(b)
+    //   sext(a + b) = sext(a) + sext(b)
     // even if the addition is not marked nsw.
     //
     // Leveraging this invarient, we can trace into an sext'ed inbound GEP
     // index if the constant offset is non-negative.
     //
     // Verified in @sext_add in split-gep.ll.
     if (ConstantInt *ConstLHS = dyn_cast<ConstantInt>(LHS)) {
       if (!ConstLHS->isNegative())
APInt ConstantOffsetExtractor::find(Value *V, bool SignExtended,
   } else if (isa<SExtInst>(V)) {
     ConstantOffset = find(U->getOperand(0), /* SignExtended */ true,
                           ZeroExtended, NonNegative).sext(BitWidth);
   } else if (isa<ZExtInst>(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.
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();
 }
 
-int64_t SeparateConstOffsetFromGEP::accumulateByteOffset(
-    GetElementPtrInst *GEP, const DataLayout *DL, bool &NeedsExtraction) {
+bool SeparateConstOffsetFromGEP::canonicalizeArrayIndicesToPointerSize(
+    GetElementPtrInst *GEP) {
+  bool Changed = false;
+  Type *IntPtrTy = DL->getIntPtrType(GEP->getType());
+  gep_type_iterator GTI = gep_type_begin(*GEP);
+  for (User::op_iterator I = GEP->op_begin() + 1, E = GEP->op_end();
+       I != E; ++I, ++GTI) {
+    // Skip struct member indices which must be i32.
+    if (isa<SequentialType>(*GTI)) {
+      if ((*I)->getType() != IntPtrTy) {
+        *I = CastInst::CreateIntegerCast(*I, IntPtrTy, true, "idxprom", GEP);
+        Changed = true;
+      }
+    }
+  }
+  return Changed;
+}
+
+bool
+SeparateConstOffsetFromGEP::convertInBoundsZExtToSExt(GetElementPtrInst *GEP) {
+  if (!GEP->isInBounds())
+    return false;
+
+  // TODO: consider alloca
+  GlobalVariable *UnderlyingObject =
+      dyn_cast<GlobalVariable>(GEP->getPointerOperand());
+  if (UnderlyingObject == nullptr)
+    return false;
+
+  uint64_t ObjectSize =
+      DL->getTypeAllocSize(UnderlyingObject->getType()->getElementType());
+  gep_type_iterator GTI = gep_type_begin(*GEP);
+  bool Changed = false;
+  for (User::op_iterator I = GEP->op_begin() + 1, E = GEP->op_end(); I != E;
+       ++I, ++GTI) {
+    if (isa<SequentialType>(*GTI)) {
+      if (ZExtInst *Extended = dyn_cast<ZExtInst>(*I)) {
+        unsigned SrcBitWidth =
+            cast<IntegerType>(Extended->getSrcTy())->getBitWidth();
+        // For GEP operand zext(a), if a <= max signed value of typeof(a), then
+        // the sign bit of a is zero and sext(a) = zext(a). Because the GEP is
+        // in bounds, we know a <= ObjectSize, so the condition can be reduced
+        // to ObjectSize <= max signed value of typeof(a).
+        if (ObjectSize <=
+            APInt::getSignedMaxValue(SrcBitWidth).getZExtValue()) {
+          *I = new SExtInst(Extended->getOperand(0), Extended->getType(),
+                            Extended->getName(), GEP);
+          Changed = true;
+        }
+      }
+    }
+  }
+  return Changed;
+}
+
+int64_t
+SeparateConstOffsetFromGEP::accumulateByteOffset(GetElementPtrInst *GEP,
+                                                 bool &NeedsExtraction) {
   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<SequentialType>(*GTI)) {
       // Tries to extract a constant offset from this GEP index.
       int64_t ConstantOffset =
           ConstantOffsetExtractor::Find(GEP->getOperand(I), DL, GEP);
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;
-  // Canonicalize array indices to pointer-size integers. This helps to simplify
-  // the logic of splitting a GEP. For example, if a + b is a pointer-size
-  // integer, we have
-  //   gep base, a + b = gep (gep base, a), b
-  // However, this equality may not hold if the size of a + b is smaller than
-  // the pointer size, because LLVM conceptually sign-extends GEP indices to
-  // pointer size before computing the address
-  // (http://llvm.org/docs/LangRef.html#id181).
-  //
-  // This canonicalization is very likely already done in clang and instcombine.
-  // Therefore, the program will probably remain the same.
-  //
-  // Verified in @i32_add in split-gep.ll
-  const DataLayout *DL = &getAnalysis<DataLayoutPass>().getDataLayout();
-  Type *IntPtrTy = DL->getIntPtrType(GEP->getType());
-  gep_type_iterator GTI = gep_type_begin(*GEP);
-  for (User::op_iterator I = GEP->op_begin() + 1, E = GEP->op_end();
-       I != E; ++I, ++GTI) {
-    if (isa<SequentialType>(*GTI)) {
-      if ((*I)->getType() != IntPtrTy) {
-        *I = CastInst::CreateIntegerCast(*I, IntPtrTy, true, "idxprom", GEP);
-        Changed = true;
-      }
-    }
-  }
+  Changed |= canonicalizeArrayIndicesToPointerSize(GEP);
+  Changed |= convertInBoundsZExtToSExt(GEP);
 
   bool NeedsExtraction;
-  int64_t AccumulativeByteOffset =
-      accumulateByteOffset(GEP, DL, NeedsExtraction);
+  int64_t AccumulativeByteOffset = accumulateByteOffset(GEP, NeedsExtraction);
 
   if (!NeedsExtraction)
     return Changed;
   // 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<TargetTransformInfo>();
   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.
-  GTI = gep_type_begin(*GEP);
+  gep_type_iterator GTI = gep_type_begin(*GEP);
   for (unsigned I = 1, E = GEP->getNumOperands(); I != E; ++I, ++GTI) {
     if (isa<SequentialType>(*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 &&
bool SeparateConstOffsetFromGEP::splitGEP(GetElementPtrInst *GEP) {
   //   %uglygep = gep %0, <offset>
   //   %new.gep = bitcast %uglygep to <type of %gep>
   //   ... %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<int64_t>(ElementTypeSizeOfGEP);
     NewGEP = GetElementPtrInst::Create(

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

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

test/Transforms/SeparateConstOffsetFromGEP/NVPTX/split-gep.ll

entry:
   %add = add nsw i32 %i, 5
   %idxprom = sext i32 %add to i64
   %p = getelementptr inbounds [1024 x %struct.S]* @struct_array, i64 0, i64 %idxprom, i32 1
   ret double* %p
 }
 ; CHECK-LABEL: @struct(
 ; CHECK: getelementptr [1024 x %struct.S]* @struct_array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i32 1
 
-; We should be able to trace into s/zext(a + b) if a + b is non-negative
+; We should be able to trace into sext(a + b) if a + b is non-negative
 ; (e.g., used as an index of an inbounds GEP) and one of a and b is
 ; non-negative.
 define float* @sext_add(i32 %i, i32 %j) {
 entry:
   %0 = add i32 %i, 1
   %1 = sext i32 %0 to i64  ; inbound sext(i + 1) = sext(i) + 1
-  %2 = sub i32 %j, 2
-  ; However, inbound sext(j - 2) != sext(j) - 2, e.g., j = INT_MIN
+  %2 = add i32 %j, -2
+  ; However, inbound sext(j + -2) != sext(j) + -2, e.g., j = INT_MIN
   %3 = sext i32 %2 to i64
   %p = getelementptr inbounds [32 x [32 x float]]* @float_2d_array, i64 0, i64 %1, i64 %3
   ret float* %p
 }
 ; CHECK-LABEL: @sext_add(
 ; CHECK-NOT: = add
+; CHECK: add i32 %j, -2
+; CHECK: sext
 ; CHECK: getelementptr [32 x [32 x float]]* @float_2d_array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}}
 ; CHECK: getelementptr float* %{{[a-zA-Z0-9]+}}, i64 32
 
 ; We should be able to trace into sext/zext if it can be distributed to both
 ; operands, e.g., sext (add nsw a, b) == add nsw (sext a), (sext b)
 ;
 ; This test verifies we can transform
 ;   gep base, a + sext(b +nsw 1), c + zext(d +nuw 1)

 ; Verifies we handle nested sext/zext correctly.
 define void @sext_zext(i32 %a, i32 %b, float** %out1, float** %out2) {
 entry:
   %0 = add nsw nuw i32 %a, 1
   %1 = sext i32 %0 to i48
   %2 = zext i48 %1 to i64    ; zext(sext(a +nsw nuw 1)) = zext(sext(a)) + 1
   %3 = add nsw i32 %b, 2
   %4 = sext i32 %3 to i48
-  %5 = zext i48 %4 to i64    ; zext(sext(a +nsw 2)) != zext(sext(a)) + 2
-  %p1 = getelementptr inbounds [32 x [32 x float]]* @float_2d_array, i64 0, i64 %2, i64 %5
+  %5 = zext i48 %4 to i64    ; zext(sext(b +nsw 2)) != zext(sext(b)) + 2
+  %p1 = getelementptr [32 x [32 x float]]* @float_2d_array, i64 0, i64 %2, i64 %5
   store float* %p1, float** %out1
   %6 = add nuw i32 %a, 3
   %7 = zext i32 %6 to i48
-  %8 = sext i48 %7 to i64 ; sext(zext(b +nuw 3)) = zext(b +nuw 3) = zext(b) + 3
+  %8 = sext i48 %7 to i64 ; sext(zext(a +nuw 3)) = zext(a +nuw 3) = zext(a) + 3
   %9 = add nsw i32 %b, 4
   %10 = zext i32 %9 to i48
   %11 = sext i48 %10 to i64  ; sext(zext(b +nsw 4)) != zext(b) + 4
-  %p2 = getelementptr inbounds [32 x [32 x float]]* @float_2d_array, i64 0, i64 %8, i64 %11
+  %p2 = getelementptr [32 x [32 x float]]* @float_2d_array, i64 0, i64 %8, i64 %11
   store float* %p2, float** %out2
   ret void
 }
 ; CHECK-LABEL: @sext_zext(
 ; CHECK: [[BASE_PTR_1:%[a-zA-Z0-9]+]] = getelementptr [32 x [32 x float]]* @float_2d_array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}}
 ; CHECK: getelementptr float* [[BASE_PTR_1]], i64 32
 ; CHECK: [[BASE_PTR_2:%[a-zA-Z0-9]+]] = getelementptr [32 x [32 x float]]* @float_2d_array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}}
 ; CHECK: getelementptr float* [[BASE_PTR_2]], i64 96
entry:
   %0 = shl i64 %a, 2
   %1 = and i64 %0, 1
   %p = getelementptr [32 x [32 x float]]* @float_2d_array, i64 0, i64 0, i64 %1
   ret float* %p
 }
 ; CHECK-LABEL: @and(
 ; CHECK: getelementptr [32 x [32 x float]]* @float_2d_array
 ; CHECK-NOT: getelementptr
+
+; if zext(a + b) <= max signed value of typeof(a + b), then we can prove
+; a + b >= 0 and zext(a + b) == sext(a + b). If we can prove further a or b is
+; non-negative, we have zext(a + b) == sext(a) + sext(b).
+define float* @inbounds_zext_add(i32 %i, i4 %j) {
+entry:
+  %0 = add i32 %i, 1
+  %1 = zext i32 %0 to i64
+  ; Because zext(i + 1) is an index of an in bounds GEP based on
+  ; float_2d_array, zext(i + 1) <= sizeof(float_2d_array) = 4096.
+  ; Furthermore, since typeof(i + 1) is i32 and 4096 < 2^31, we are sure the
+  ; sign bit of i + 1 is 0. This implies zext(i + 1) = sext(i + 1).
+  %2 = add i4 %j, 2
+  %3 = zext i4 %2 to i64
+  ; In this case, typeof(j + 2) is i4, so zext(j + 2) <= 4096 does not imply
+  ; the sign bit of j + 2 is 0.
+  %p = getelementptr inbounds [32 x [32 x float]]* @float_2d_array, i64 0, i64 %1, i64 %3
+  ret float* %p
+}
+; CHECK-LABEL: @inbounds_zext_add(
+; CHECK-NOT: add
+; CHECK: add i4 %j, 2
+; CHECK: sext
+; CHECK: getelementptr [32 x [32 x float]]* @float_2d_array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}}
+; CHECK: getelementptr float* %{{[a-zA-Z0-9]+}}, i64 32