[llvm] [GISel] Set more MIFlags when translating GEPs (PR #151708)

[Fabian Ritter via llvm-commits]

================
@@ -1592,9 +1592,19 @@ bool IRTranslator::translateGetElementPtr(const User &U,
   Type *OffsetIRTy = DL->getIndexType(PtrIRTy);
   LLT OffsetTy = getLLTForType(*OffsetIRTy, *DL);
 
-  uint32_t Flags = 0;
+  uint32_t PtrAddFlags = 0;
+  // Each PtrAdd generated to implement the GEP inherits its nuw, nusw, inbounds
+  // flags.
   if (const Instruction *I = dyn_cast<Instruction>(&U))
-    Flags = MachineInstr::copyFlagsFromInstruction(*I);
+    PtrAddFlags = MachineInstr::copyFlagsFromInstruction(*I);
+
+  auto PtrAddFlagsWithConst = [&](int64_t Offset) {
+    // For nusw/inbounds GEP with an offset that is nonnegative when interpreted
+    // as signed, assume there is no unsigned overflow.
+    if (Offset >= 0 && (PtrAddFlags & MachineInstr::MIFlag::NoUSWrap))
+      return PtrAddFlags | MachineInstr::MIFlag::NoUWrap;
+    return PtrAddFlags;
+  };
----------------
ritter-x2a wrote:

The middle-end doesn't infer it in all cases, as far as I can tell: When I compile the following HIP code:

```
struct S {
  int a;
  int b[3];
  int c;
};

__global__ void foo(S *p, int offset) {
  p[offset].b[2] = 42;
}
```

with
```
clang -xhip --offload-arch=gfx1030 --driver-mode=g++ --offload-device-only -fglobal-isel -O3 -mllvm -print-before-all test.hip
```

I get this IR, without NUW, before the IRTranslator:
```
  ...
  %0 = getelementptr inbounds %struct.S, ptr addrspace(1) %p.coerce.load, i64 %idxprom, i32 1, i64 0
  %arrayidx23 = getelementptr inbounds i8, ptr addrspace(1) %0, i64 8
  ...
```
Without this check here, the `G_PTR_ADD`s in the generated MIR don't have NUW, with it the `G_PTR_ADD`s with a constant operand do.

I think part of the reason for that is that the last instance of instcombine runs before separate-const-offset-from-gep, so the [relevant transform in instcombine](https://github.com/llvm/llvm-project/blob/1862e3c56ce1eb3762b4ccd6e9e222f3022c9c7c/llvm/lib/Transforms/InstCombine/InstructionCombining.cpp#L3409-L3416) only sees a single multi-offset GEP with one potentially negative offset, so it can't set NUW for the entire GEP.

https://github.com/llvm/llvm-project/pull/151708