[InferAddressSpaces] Handle the pair of `ptrtoint`/`inttoptr`.

[hliao]

Summary

• ptrtoint and inttoptr are defined as no-op casts if the integer value as the same size as the pointer value. The pair of ptrtoint/inttoptr is in fact a no-op cast sequence between different address spaces. Teach infer-address-spaces to handle them like a bitcast.

SROA part is split into https://reviews.llvm.org/D81943

[lebedev.ri]

This should be two separate patches - inferaddressspace and SROA.

[arsenm]

Should be separate patches for each pass

llvm/test/Transforms/InferAddressSpaces/noop-ptrint-pair.ll

17

This is missing a number of test cases I would like to see. First, I would like to make sure if the intermediate size is larger or smaller than the pointer size, this isn't done. I also want to stress the full combination of cases in a constant expression, not just instructions.

[hliao]

In D81938#2095828, @lebedev.ri wrote:

This should be two separate patches - inferaddressspace and SROA.

Yes, I prepared that into 2 commits but arc combines them together.

[tra]

In D81938#2095869, @hliao wrote:

In D81938#2095828, @lebedev.ri wrote:

This should be two separate patches - inferaddressspace and SROA.

Yes, I prepared that into 2 commits but arc combines them together.

I usually put the patches into separate branches, with one having the other one as an upstream branch and then do arc diff with each branch checked out.
The downside is that you will need to rebase infer-branch every time you update sroa-branch.
Another approach that may work is to set arc config to only consider the checked out commit only. arc set-config base "arc:this, arc:prompt" should do that.

[hliao]

Fix constant expression handling.

[hliao]

In D81938#2095982, @tra wrote:

In D81938#2095869, @hliao wrote:

In D81938#2095828, @lebedev.ri wrote:

This should be two separate patches - inferaddressspace and SROA.

Yes, I prepared that into 2 commits but arc combines them together.

I usually put the patches into separate branches, with one having the other one as an upstream branch and then do arc diff with each branch checked out.
The downside is that you will need to rebase infer-branch every time you update sroa-branch.
Another approach that may work is to set arc config to only consider the checked out commit only. arc set-config base "arc:this, arc:prompt" should do that.

Thanks for the ar option approach. I usually do the same thing creating a bunch of local branches to separate individual changes. It's quite cumbersome to maintain them. More than 80 local branches are really difficult to maintain. Will try that arc option.

[arsenm]

I'm not entirely convinced this is safe in all contexts. I think you can argue that this is safe if it directly feeds a memory instruction, as the access would be undefined if it weren't valid to do the no-op cast. However, I'm not sure if this is safe if used purely in arithmetic contexts. If you're just comparing the reinterpreted pointer values for example, I don't think that would be undefined

[hliao]

In D81938#2096500, @arsenm wrote:

I'm not entirely convinced this is safe in all contexts. I think you can argue that this is safe if it directly feeds a memory instruction, as the access would be undefined if it weren't valid to do the no-op cast. However, I'm not sure if this is safe if used purely in arithmetic contexts. If you're just comparing the reinterpreted pointer values for example, I don't think that would be undefined

Would it be safe if we double-check the target hook and ensure that's a no-op addrspacecast between address spaces?

[arsenm]

In D81938#2096732, @hliao wrote:

In D81938#2096500, @arsenm wrote:

I'm not entirely convinced this is safe in all contexts. I think you can argue that this is safe if it directly feeds a memory instruction, as the access would be undefined if it weren't valid to do the no-op cast. However, I'm not sure if this is safe if used purely in arithmetic contexts. If you're just comparing the reinterpreted pointer values for example, I don't think that would be undefined

Would it be safe if we double-check the target hook and ensure that's a no-op addrspacecast between address spaces?

Yes, that would be safer since you can see there's no difference

[hliao]

Add TTI hooks to double-check that address space casting is no-op.

NOTE: This change depends on - https://reviews.llvm.org/D82025 (exposing isNoopAddrSpaceCast from TLI into TTI) to be compiled. - https://reviews.llvm.org/D81943 (enhance SROA to handle no-op address space casting) to pass all testing.

[hliao]

PING

[hliao]

Rebase to trunk. All prerequisites are landed.

[hliao]

ping for code review

[arsenm]

clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu

62–64

Positive checks are preferable (here and through the rest of the file)

llvm/lib/Transforms/Scalar/InferAddressSpaces.cpp

223

Can you add a comment elaborating on why this is necessary? This is special casing this to paper over the missing no-op pointer bitcast

237–239

Can you also elaborate on why the isNoopAddrSpaceCast check is needed? I'm not 100% sure it is in all contexts, so want to be sure to document that

241

data layout should be const reference

llvm/test/Transforms/InferAddressSpaces/noop-ptrint-pair.ll

22

Can you also add one where the cast would be a no-op, but the integer size is different? e.g. 0 to 1 and 1 to 0 where the intermediate size is 32-bits or 128 bits

46

Can you add some cases where the pointer value isn't used for memory? Like pure pointer arithmetic

[hliao]

clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu

62–64

here, we need to ensure, after optimization, these values are promoted from memory

[hliao]

Add comments and enhance tests.

[hliao]

llvm/lib/Transforms/Scalar/InferAddressSpaces.cpp

237–239

As we may convert a pair of ptr/int casts into an addrspacecast if both of them are no-op casts, if that's not a no-op addrspacecast under a specific target, we may introduce an invalid addrspacecast based on the current IR spec.

[hliao]

enhance test more.

[arsenm]

clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu

62–64

Yes, you can explicitly check for the instructions you expect, rather than ones you do not

llvm/lib/Transforms/Scalar/InferAddressSpaces.cpp

237–239

I mean there needs to be comment explaining all of this

[hliao]

Revise again.

[hliao]

Rebase to the trunk.

[hliao]

ping for code review

[arsenm]

clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu

62

Probably should use filecheck variables to avoid depending on the numbers

65

I think relying on the name may fail in a release build, filecheck variable

llvm/lib/Transforms/Scalar/InferAddressSpaces.cpp

237–239

I think this is missing one of the key points I'm worried about for the IR semantics.

I believe if you were to invalidly reinterpret a pointer with another address space, it would be undefined to access the invalid pointer and thus you can use this to infer. In pure arithmetic cases, I'm not sure this would hold. This check is to avoid breaking cases where it may still be valid to do something with reinterpreted pointer bits other than access memory

llvm/test/Transforms/InferAddressSpaces/noop-ptrint-pair.ll

3

The pass early exits without TTI since there's no flat address space value, so this isn't really testing anything. We would have to add a flag for the flat value to bypass it.

[hliao]

Revise

[arsenm]

llvm/lib/Transforms/Scalar/InferAddressSpaces.cpp

232–233

The grammar for the first sentence is off. "Check if the pair of ptrtoint/inttoptr is a no-op cast."

240

No "an".

"if that pionter is dereferenced" would be more specific than than "pointer bits accessed"

[hliao]

Revie the grammar. s/is/as/ in the first sentence.

[hliao]

remove 'an'

clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu

-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck --check-prefixes=COMMON,CHECK %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -disable-O0-optnone -o - | opt -S -O2 | FileCheck %s --check-prefixes=COMMON,OPT
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -x hip %s -o - | FileCheck -check-prefix=HOST %s
 
 #include "Inputs/cuda.h"
 
 // Coerced struct from `struct S` without all generic pointers lowered into
 // global ones.
-// CHECK: %struct.S.coerce = type { i32 addrspace(1)*, float addrspace(1)* }
-// CHECK: %struct.T.coerce = type { [2 x float addrspace(1)*] }
+// COMMON: %struct.S.coerce = type { i32 addrspace(1)*, float addrspace(1)* }
+// COMMON: %struct.T.coerce = type { [2 x float addrspace(1)*] }
 
 // On the host-side compilation, generic pointer won't be coerced.
 // HOST-NOT: %struct.S.coerce
 // HOST-NOT: %struct.T.coerce
 
-// CHECK: define amdgpu_kernel void  @_Z7kernel1Pi(i32 addrspace(1)* %x.coerce)
 // HOST: define void @_Z22__device_stub__kernel1Pi(i32* %x)
+// COMMON-LABEL: define amdgpu_kernel void @_Z7kernel1Pi(i32 addrspace(1)*{{.*}} %x.coerce)
+// CHECK:     = addrspacecast [[TYPE:.*]] addrspace(1)* %{{.*}} to [[TYPE]]*
+// CHECK-NOT: = addrspacecast [[TYPE:.*]] addrspace(1)* %{{.*}} to [[TYPE]]*
+// OPT-NOT: alloca
+// OPT-NOT: inttoptr
 __global__ void kernel1(int *x) {
   x[0]++;
 }
 
-// CHECK: define amdgpu_kernel void  @_Z7kernel2Ri(i32 addrspace(1)* nonnull align 4 dereferenceable(4) %x.coerce)
 // HOST: define void @_Z22__device_stub__kernel2Ri(i32* nonnull align 4 dereferenceable(4) %x)
+// COMMON-LABEL: define amdgpu_kernel void @_Z7kernel2Ri(i32 addrspace(1)*{{.*}} nonnull align 4 dereferenceable(4) %x.coerce)
+// CHECK:     = addrspacecast [[TYPE:.*]] addrspace(1)* %{{.*}} to [[TYPE]]*
+// CHECK-NOT: = addrspacecast [[TYPE:.*]] addrspace(1)* %{{.*}} to [[TYPE]]*
+// OPT-NOT: alloca
+// OPT-NOT: ptrtoint
+// OPT-NOT: inttoptr
 __global__ void kernel2(int &x) {
   x++;
 }
 
-// CHECK: define amdgpu_kernel void  @_Z7kernel3PU3AS2iPU3AS1i(i32 addrspace(2)* %x, i32 addrspace(1)* %y)
 // HOST: define void @_Z22__device_stub__kernel3PU3AS2iPU3AS1i(i32 addrspace(2)* %x, i32 addrspace(1)* %y)
+// CHECK-LABEL: define amdgpu_kernel void  @_Z7kernel3PU3AS2iPU3AS1i(i32 addrspace(2)*{{.*}} %x, i32 addrspace(1)*{{.*}} %y)
+// CHECK-NOT: = addrspacecast [[TYPE:.*]] addrspace(1)* %{{.*}} to [[TYPE]]*
 __global__ void kernel3(__attribute__((address_space(2))) int *x,
                         __attribute__((address_space(1))) int *y) {
   y[0] = x[0];
 }
 
-// CHECK: define void @_Z4funcPi(i32* %x)
+// COMMON-LABEL: define void @_Z4funcPi(i32*{{.*}} %x)
+// CHECK-NOT: = addrspacecast [[TYPE:.*]] addrspace(1)* %{{.*}} to [[TYPE]]*
+// OPT-NOT: alloca
+// OPT-NOT: ptrtoint
+// OPT-NOT: inttoptr
 __device__ void func(int *x) {
   x[0]++;
 }
 
 struct S {
   int *x;
   float *y;
 };
 // `by-val` struct will be coerced into a similar struct with all generic
 // pointers lowerd into global ones.
-// CHECK: define amdgpu_kernel void @_Z7kernel41S(%struct.S.coerce %s.coerce)
 // HOST: define void @_Z22__device_stub__kernel41S(i32* %s.coerce0, float* %s.coerce1)
+// COMMON-LABEL: define amdgpu_kernel void @_Z7kernel41S(%struct.S.coerce %s.coerce)
+// OPT-NOT: alloca
+// OPT-NOT: ptrtoint
+// OPT-NOT: inttoptr
 __global__ void kernel4(struct S s) {
   s.x[0]++;
   s.y[0] += 1.f;
 }
 
 // If a pointer to struct is passed, only the pointer itself is coerced into the global one.
-// CHECK: define amdgpu_kernel void @_Z7kernel5P1S(%struct.S addrspace(1)* %s.coerce)
 // HOST: define void @_Z22__device_stub__kernel5P1S(%struct.S* %s)
+// COMMON-LABEL: define amdgpu_kernel void @_Z7kernel5P1S(%struct.S addrspace(1)*{{.*}} %s.coerce)
+// OPT-NOT: alloca
+// OPT-NOT: ptrtoint
+// OPT-NOT: inttoptr
 __global__ void kernel5(struct S *s) {
   s->x[0]++;
   s->y[0] += 1.f;
 }
 
 struct T {
   float *x[2];
 };
 // `by-val` array is also coerced.
-// CHECK: define amdgpu_kernel void @_Z7kernel61T(%struct.T.coerce %t.coerce)
 // HOST: define void @_Z22__device_stub__kernel61T(float* %t.coerce0, float* %t.coerce1)
+// COMMON-LABEL: define amdgpu_kernel void @_Z7kernel61T(%struct.T.coerce %t.coerce)
+// OPT-NOT: alloca
+// OPT-NOT: ptrtoint
+// OPT-NOT: inttoptr
 __global__ void kernel6(struct T t) {
   t.x[0][0] += 1.f;
   t.x[1][0] += 2.f;
 }
 
 // Check that coerced pointers retain the noalias attribute when qualified with __restrict.
-// CHECK: define amdgpu_kernel void @_Z7kernel7Pi(i32 addrspace(1)* noalias %x.coerce)
 // HOST: define void @_Z22__device_stub__kernel7Pi(i32* noalias %x)
+// COMMON-LABEL: define amdgpu_kernel void @_Z7kernel7Pi(i32 addrspace(1)* noalias{{.*}} %x.coerce)
+// OPT-NOT: alloca
+// OPT-NOT: ptrtoint
+// OPT-NOT: inttoptr
 __global__ void kernel7(int *__restrict x) {
   x[0]++;
 }

llvm/lib/Transforms/Scalar/InferAddressSpaces.cpp

 namespace {
 
 using ValueToAddrSpaceMapTy = DenseMap<const Value *, unsigned>;
 using PostorderStackTy = llvm::SmallVector<PointerIntPair<Value *, 1, bool>, 4>;
 
 /// InferAddressSpaces
 class InferAddressSpaces : public FunctionPass {
   const TargetTransformInfo *TTI = nullptr;
+  const DataLayout *DL = nullptr;
 
   /// Target specific address space which uses of should be replaced if
   /// possible.
   unsigned FlatAddrSpace = 0;
 
 public:
   static char ID;
 

 
 void initializeInferAddressSpacesPass(PassRegistry &);
 
 } // end namespace llvm
 
 INITIALIZE_PASS(InferAddressSpaces, DEBUG_TYPE, "Infer address spaces",
                 false, false)
 
+static bool isNoopPtrIntCastPair(const Operator *I2P, const DataLayout *DL) {
+  assert(I2P->getOpcode() == Instruction::IntToPtr);
+  auto *P2I = dyn_cast<Operator>(I2P->getOperand(0));
+  if (!P2I || P2I->getOpcode() != Instruction::PtrToInt)
+    return false;
+  // The pair of `ptrtoint`/`inttoptr` is a no-op cast if both of them are
+  // no-op casts.
+  return CastInst::isNoopCast(Instruction::CastOps(I2P->getOpcode()),
+                              I2P->getOperand(0)->getType(), I2P->getType(),
+                              *DL) &&
+         CastInst::isNoopCast(Instruction::CastOps(P2I->getOpcode()),
+                              P2I->getOperand(0)->getType(), P2I->getType(),
+                              *DL);
+}
+
 // Returns true if V is an address expression.
 // TODO: Currently, we consider only phi, bitcast, addrspacecast, and
 // getelementptr operators.
-static bool isAddressExpression(const Value &V) {
+static bool isAddressExpression(const Value &V, const DataLayout *DL) {
   const Operator *Op = dyn_cast<Operator>(&V);
   if (!Op)
     return false;
 
   switch (Op->getOpcode()) {
   case Instruction::PHI:
     assert(Op->getType()->isPointerTy());
     return true;
   case Instruction::BitCast:
   case Instruction::AddrSpaceCast:
   case Instruction::GetElementPtr:
     return true;
   case Instruction::Select:
     return Op->getType()->isPointerTy();
   case Instruction::Call: {
     const IntrinsicInst *II = dyn_cast<IntrinsicInst>(&V);
     return II && II->getIntrinsicID() == Intrinsic::ptrmask;
   }
+  case Instruction::IntToPtr:
+    return isNoopPtrIntCastPair(Op, DL);
   default:
     return false;
   }
 }
 
 // Returns the pointer operands of V.
 //
 // Precondition: V is an address expression.
-static SmallVector<Value *, 2> getPointerOperands(const Value &V) {
+static SmallVector<Value *, 2> getPointerOperands(const Value &V,
+                                                  const DataLayout *DL) {
   const Operator &Op = cast<Operator>(V);
   switch (Op.getOpcode()) {
   case Instruction::PHI: {
     auto IncomingValues = cast<PHINode>(Op).incoming_values();
     return SmallVector<Value *, 2>(IncomingValues.begin(),
                                    IncomingValues.end());
   }
   case Instruction::BitCast:
   case Instruction::AddrSpaceCast:
   case Instruction::GetElementPtr:
     return {Op.getOperand(0)};
   case Instruction::Select:
     return {Op.getOperand(1), Op.getOperand(2)};
   case Instruction::Call: {
     const IntrinsicInst &II = cast<IntrinsicInst>(Op);
     assert(II.getIntrinsicID() == Intrinsic::ptrmask &&
            "unexpected intrinsic call");
     return {II.getArgOperand(0)};
   }
+  case Instruction::IntToPtr: {
+    assert(isNoopPtrIntCastPair(&Op, DL));
+    auto *P2I = cast<Operator>(Op.getOperand(0));
+    return {P2I->getOperand(0)};
+  }
   default:
     llvm_unreachable("Unexpected instruction type.");
   }
 }
 
 bool InferAddressSpaces::rewriteIntrinsicOperands(IntrinsicInst *II,
                                                   Value *OldV,
                                                   Value *NewV) const {
void InferAddressSpaces::appendsFlatAddressExpressionToPostorderStack(
     Value *V, PostorderStackTy &PostorderStack,
     DenseSet<Value *> &Visited) const {
   assert(V->getType()->isPointerTy());
 
   // Generic addressing expressions may be hidden in nested constant
   // expressions.
   if (ConstantExpr *CE = dyn_cast<ConstantExpr>(V)) {
     // TODO: Look in non-address parts, like icmp operands.
-    if (isAddressExpression(*CE) && Visited.insert(CE).second)
+    if (isAddressExpression(*CE, DL) && Visited.insert(CE).second)
       PostorderStack.emplace_back(CE, false);
 
     return;
   }
 
-  if (isAddressExpression(*V) &&
+  if (isAddressExpression(*V, DL) &&
       V->getType()->getPointerAddressSpace() == FlatAddrSpace) {
     if (Visited.insert(V).second) {
       PostorderStack.emplace_back(V, false);
 
       Operator *Op = cast<Operator>(V);
       for (unsigned I = 0, E = Op->getNumOperands(); I != E; ++I) {
         if (ConstantExpr *CE = dyn_cast<ConstantExpr>(Op->getOperand(I))) {
-          if (isAddressExpression(*CE) && Visited.insert(CE).second)
+          if (isAddressExpression(*CE, DL) && Visited.insert(CE).second)
             PostorderStack.emplace_back(CE, false);
         }
       }
     }
   }
 }
 
 // Returns all flat address expressions in function F. The elements are ordered
else if (ICmpInst *Cmp = dyn_cast<ICmpInst>(&I)) {
       // FIXME: Handle vectors of pointers
       if (Cmp->getOperand(0)->getType()->isPointerTy()) {
         PushPtrOperand(Cmp->getOperand(0));
         PushPtrOperand(Cmp->getOperand(1));
       }
     } else if (auto *ASC = dyn_cast<AddrSpaceCastInst>(&I)) {
       if (!ASC->getType()->isVectorTy())
         PushPtrOperand(ASC->getPointerOperand());
+    } else if (auto *I2P = dyn_cast<IntToPtrInst>(&I)) {
+      if (isNoopPtrIntCastPair(cast<Operator>(I2P), DL))
+        PushPtrOperand(
+            cast<PtrToIntInst>(I2P->getOperand(0))->getPointerOperand());
     }
   }
 
   std::vector<WeakTrackingVH> Postorder; // The resultant postorder.
   while (!PostorderStack.empty()) {
     Value *TopVal = PostorderStack.back().getPointer();
     // If the operands of the expression on the top are already explored,
     // adds that expression to the resultant postorder.
     if (PostorderStack.back().getInt()) {
       if (TopVal->getType()->getPointerAddressSpace() == FlatAddrSpace)
         Postorder.push_back(TopVal);
       PostorderStack.pop_back();
       continue;
     }
     // Otherwise, adds its operands to the stack and explores them.
     PostorderStack.back().setInt(true);
-    for (Value *PtrOperand : getPointerOperands(*TopVal)) {
+    for (Value *PtrOperand : getPointerOperands(*TopVal, DL)) {
       appendsFlatAddressExpressionToPostorderStack(PtrOperand, PostorderStack,
                                                    Visited);
     }
   }
   return Postorder;
 }
 
 // A helper function for cloneInstructionWithNewAddressSpace. Returns the clone
GetElementPtrInst *NewGEP = GetElementPtrInst::Create(
         SmallVector<Value *, 4>(GEP->idx_begin(), GEP->idx_end()));
     NewGEP->setIsInBounds(GEP->isInBounds());
     return NewGEP;
   }
   case Instruction::Select:
     assert(I->getType()->isPointerTy());
     return SelectInst::Create(I->getOperand(0), NewPointerOperands[1],
                               NewPointerOperands[2], "", nullptr, I);
+  case Instruction::IntToPtr: {
+    assert(isNoopPtrIntCastPair(cast<Operator>(I), DL));
+    Value *Src = cast<Operator>(I->getOperand(0))->getOperand(0);
+    assert(Src->getType()->getPointerAddressSpace() == NewAddrSpace);
+    if (Src->getType() != NewPtrType)
+      return new BitCastInst(Src, NewPtrType);
+    return Src;
+  }
   default:
     llvm_unreachable("Unexpected opcode");
   }
 }
 
 // Similar to cloneInstructionWithNewAddressSpace, returns a clone of the
 // constant expression `CE` with its operands replaced as specified in
 // ValueWithNewAddrSpace.

 // expression whose address space needs to be modified, in postorder.
 //
 // See cloneInstructionWithNewAddressSpace for the meaning of UndefUsesToFix.
 Value *InferAddressSpaces::cloneValueWithNewAddressSpace(
   Value *V, unsigned NewAddrSpace,
   const ValueToValueMapTy &ValueWithNewAddrSpace,
   SmallVectorImpl<const Use *> *UndefUsesToFix) const {
   // All values in Postorder are flat address expressions.
-  assert(isAddressExpression(*V) &&
+  assert(isAddressExpression(*V, DL) &&
          V->getType()->getPointerAddressSpace() == FlatAddrSpace);
 
   if (Instruction *I = dyn_cast<Instruction>(V)) {
     Value *NewV = cloneInstructionWithNewAddressSpace(
       I, NewAddrSpace, ValueWithNewAddrSpace, UndefUsesToFix);
     if (Instruction *NewI = dyn_cast_or_null<Instruction>(NewV)) {
       if (NewI->getParent() == nullptr) {
         NewI->insertBefore(I);
unsigned InferAddressSpaces::joinAddressSpaces(unsigned AS1,
   return (AS1 == AS2) ? AS1 : FlatAddrSpace;
 }
 
 bool InferAddressSpaces::runOnFunction(Function &F) {
   if (skipFunction(F))
     return false;
 
   TTI = &getAnalysis<TargetTransformInfoWrapperPass>().getTTI(F);
+  DL = &F.getParent()->getDataLayout();
 
   if (FlatAddrSpace == UninitializedAddressSpace) {
     FlatAddrSpace = TTI->getFlatAddressSpace();
     if (FlatAddrSpace == UninitializedAddressSpace)
       return false;
   }
 
   // Collects all flat address expressions in postorder.
if (Op.getOpcode() == Instruction::Select) {
 
     if (C0 && isSafeToCastConstAddrSpace(C0, Src1AS))
       NewAS = Src1AS;
     else if (C1 && isSafeToCastConstAddrSpace(C1, Src0AS))
       NewAS = Src0AS;
     else
       NewAS = joinAddressSpaces(Src0AS, Src1AS);
   } else {
-    for (Value *PtrOperand : getPointerOperands(V)) {
+    for (Value *PtrOperand : getPointerOperands(V, DL)) {
       auto I = InferredAddrSpace.find(PtrOperand);
       unsigned OperandAS = I != InferredAddrSpace.end() ?
         I->second : PtrOperand->getType()->getPointerAddressSpace();
 
       // join(flat, *) = flat. So we can break if NewAS is already flat.
       NewAS = joinAddressSpaces(NewAS, OperandAS);
       if (NewAS == FlatAddrSpace)
         break;

llvm/lib/Transforms/Scalar/SROA.cpp

if (!NewTy->isSingleValueType() || !OldTy->isSingleValueType())
     return false;
 
   // We can convert pointers to integers and vice-versa. Same for vectors
   // of pointers and integers.
   OldTy = OldTy->getScalarType();
   NewTy = NewTy->getScalarType();
   if (NewTy->isPointerTy() || OldTy->isPointerTy()) {
     if (NewTy->isPointerTy() && OldTy->isPointerTy()) {
-      return cast<PointerType>(NewTy)->getPointerAddressSpace() ==
-        cast<PointerType>(OldTy)->getPointerAddressSpace();
+      unsigned OldAS = cast<PointerType>(OldTy)->getPointerAddressSpace();
+      unsigned NewAS = cast<PointerType>(NewTy)->getPointerAddressSpace();
+      return OldAS == NewAS ||
+             DL.getPointerSize(OldAS) == DL.getPointerSize(NewAS);
     }
 
     // We can convert integers to integral pointers, but not to non-integral
     // pointers.
     if (OldTy->isIntegerTy())
       return !DL.isNonIntegralPointerType(NewTy);
 
     // We can convert integral pointers to integers, but non-integral pointers
if (OldTy->isPtrOrPtrVectorTy() && NewTy->isIntOrIntVectorTy()) {
     // Expand i8* to <2 x i32> --> i8* to i64 to <2 x i32>
     if (!OldTy->isVectorTy() && NewTy->isVectorTy())
       return IRB.CreateBitCast(IRB.CreatePtrToInt(V, DL.getIntPtrType(OldTy)),
                                NewTy);
 
     return IRB.CreatePtrToInt(V, NewTy);
   }
 
+  if (OldTy->isPtrOrPtrVectorTy() && NewTy->isPtrOrPtrVectorTy()) {
+    unsigned OldAS = OldTy->getPointerAddressSpace();
+    unsigned NewAS = NewTy->getPointerAddressSpace();
+    // Generate pair of ptrtoint/inttoptr.
+    if (OldAS != NewAS) {
+      assert(DL.getPointerSize(OldAS) == DL.getPointerSize(NewAS));
+      return IRB.CreateIntToPtr(IRB.CreatePtrToInt(V, DL.getIntPtrType(OldTy)),
+                                NewTy);
+    }
+  }
+
   return IRB.CreateBitCast(V, NewTy);
 }
 
 /// Test whether the given slice use can be promoted to a vector.
 ///
 /// This function is called to test each entry in a partition which is slated
 /// for a single slice.
 static bool isVectorPromotionViableForSlice(Partition &P, const Slice &S,

llvm/test/Transforms/InferAddressSpaces/noop-ptrint-pair.ll

+; RUN: opt -S -o - -sroa -infer-address-spaces %s | FileCheck %s
+
+target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-ni:7"
+target triple = "amdgcn-amd-amdhsa"
+
+; CHECK-LABEL: @noop_ptrint_pair(
+; CHECK-NEXT: store i32 0, i32 addrspace(1)* %{{.*}}
+; CHECK-NEXT: ret void
+define void @noop_ptrint_pair(i32 addrspace(1)* %x.coerce) {
+  %x = alloca i32*, align 8, addrspace(5)
+  %x1 = addrspacecast i32* addrspace(5)* %x to i32**
+  %x2 = bitcast i32** %x1 to i32 addrspace(1)**
+  store i32 addrspace(1)* %x.coerce, i32 addrspace(1)** %x2
+  %x3 = load i32*, i32** %x1
+  store i32 0, i32* %x3
+  ret void
+}

llvm/test/Transforms/SROA/address-spaces.ll

 ; RUN: opt < %s -sroa -S | FileCheck %s
-target datalayout = "e-p:64:64:64-p1:16:16:16-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:32:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-n8:16:32:64"
+target datalayout = "e-p:64:64:64-p1:16:16:16-p3:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:32:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-n8:16:32:64"
 
 declare void @llvm.memcpy.p0i8.p0i8.i32(i8* nocapture, i8* nocapture readonly, i32, i1)
 declare void @llvm.memcpy.p1i8.p0i8.i32(i8 addrspace(1)* nocapture, i8* nocapture readonly, i32, i1)
 declare void @llvm.memcpy.p0i8.p1i8.i32(i8* nocapture, i8 addrspace(1)* nocapture readonly, i32, i1)
 declare void @llvm.memcpy.p1i8.p1i8.i32(i8 addrspace(1)* nocapture, i8 addrspace(1)* nocapture readonly, i32, i1)
 
 
 ; Make sure an illegal bitcast isn't introduced
for.end:
   ret void
 }
  
 %union.anon = type { i32* }
 
 @g = common global i32 0, align 4
 @l = common addrspace(3) global i32 0, align 4
 
-; Make sure an illegal bitcast isn't introduced
+; If pointers from different address spaces have different sizes, make sure an
+; illegal bitcast isn't introduced
 define void @pr27557() {
 ; CHECK-LABEL: @pr27557(
 ; CHECK: %[[CAST:.*]] = bitcast i32** {{.*}} to i32 addrspace(3)**
 ; CHECK: store i32 addrspace(3)* @l, i32 addrspace(3)** %[[CAST]]
   %1 = alloca %union.anon, align 8
   %2 = bitcast %union.anon* %1 to i32**
   store i32* @g, i32** %2, align 8
   %3 = bitcast %union.anon* %1 to i32 addrspace(3)**
   store i32 addrspace(3)* @l, i32 addrspace(3)** %3, align 8
   ret void
 }
 
+@l2 = common addrspace(2) global i32 0, align 4
+
+; If pointers from different address spaces have the same size, that pointer
+; should be promoted through the pair of `ptrtoint`/`inttoptr`.
+define i32* @pr27557.alt() {
+; CHECK-LABEL: @pr27557.alt(
+; CHECK: ret i32* inttoptr (i64 ptrtoint (i32 addrspace(2)* @l2 to i64) to i32*)
+  %1 = alloca %union.anon, align 8
+  %2 = bitcast %union.anon* %1 to i32 addrspace(2)**
+  store i32 addrspace(2)* @l2, i32 addrspace(2)** %2, align 8
+  %3 = bitcast %union.anon* %1 to i32**
+  %4 = load i32*, i32** %3, align 8
+  ret i32* %4
+}
+
 ; Make sure pre-splitting doesn't try to introduce an illegal bitcast
 define float @presplit(i64 addrspace(1)* %p) {
 entry:
 ; CHECK-LABEL: @presplit(
 ; CHECK: %[[CAST:.*]] = bitcast i64 addrspace(1)* {{.*}} to i32 addrspace(1)*
 ; CHECK: load i32, i32 addrspace(1)* %[[CAST]]
    %b = alloca i64
    %b.cast = bitcast i64* %b to [2 x float]*

llvm/test/Transforms/SROA/alloca-address-space.ll

 ; RUN: opt < %s -sroa -S | FileCheck %s
-target datalayout = "e-p:64:64:64-p1:16:16:16-p2:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:32:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-n8:16:32:64-A2"
+target datalayout = "e-p:64:64:64-p1:16:16:16-p2:32:32-p3:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:32:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-n8:16:32:64-A2"
 
 declare void @llvm.memcpy.p2i8.p2i8.i32(i8 addrspace(2)* nocapture, i8 addrspace(2)* nocapture readonly, i32, i1)
 declare void @llvm.memcpy.p1i8.p2i8.i32(i8 addrspace(1)* nocapture, i8 addrspace(2)* nocapture readonly, i32, i1)
 declare void @llvm.memcpy.p2i8.p1i8.i32(i8 addrspace(2)* nocapture, i8 addrspace(1)* nocapture readonly, i32, i1)
 declare void @llvm.memcpy.p1i8.p1i8.i32(i8 addrspace(1)* nocapture, i8 addrspace(1)* nocapture readonly, i32, i1)
 
 
 
for.end:
   ret void
 }
 
 %union.anon = type { i32* }
 
 @g = common global i32 0, align 4
 @l = common addrspace(3) global i32 0, align 4
 
-; Make sure an illegal bitcast isn't introduced
+; If pointers from different address spaces have different sizes, make sure an
+; illegal bitcast isn't introduced
 ; CHECK-LABEL: @pr27557(
 ; CHECK: %[[CAST:.*]] = bitcast i32* addrspace(2)* {{.*}} to i32 addrspace(3)* addrspace(2)*
 ; CHECK: store i32 addrspace(3)* @l, i32 addrspace(3)* addrspace(2)* %[[CAST]]
 define void @pr27557() {
   %1 = alloca %union.anon, align 8, addrspace(2)
   %2 = bitcast %union.anon addrspace(2)* %1 to i32* addrspace(2)*
   store i32* @g, i32* addrspace(2)* %2, align 8
   %3 = bitcast %union.anon addrspace(2)* %1 to i32 addrspace(3)* addrspace(2)*
   store i32 addrspace(3)* @l, i32 addrspace(3)* addrspace(2)* %3, align 8
   ret void
 }
 
+@l4 = common addrspace(4) global i32 0, align 4
+
+; If pointers from different address spaces have the same size, that pointer
+; should be promoted through the pair of `ptrtoint`/`inttoptr`.
+define i32* @pr27557.alt() {
+; CHECK-LABEL: @pr27557.alt(
+; CHECK: ret i32* inttoptr (i64 ptrtoint (i32 addrspace(4)* @l4 to i64) to i32*)
+  %1 = alloca %union.anon, align 8, addrspace(2)
+  %2 = bitcast %union.anon addrspace(2)* %1 to i32 addrspace(4)* addrspace(2)*
+  store i32 addrspace(4)* @l4, i32 addrspace(4)* addrspace(2)* %2, align 8
+  %3 = bitcast %union.anon addrspace(2)* %1 to i32* addrspace(2)*
+  %4 = load i32*, i32* addrspace(2)* %3, align 8
+  ret i32* %4
+}
+
 ; Test load from and store to non-zero address space.
 define void @test_load_store_diff_addr_space([2 x float] addrspace(1)* %complex1, [2 x float] addrspace(1)* %complex2) {
 ; CHECK-LABEL: @test_load_store_diff_addr_space
 ; CHECK-NOT: alloca
 ; CHECK: load i32, i32 addrspace(1)*
 ; CHECK: load i32, i32 addrspace(1)*
 ; CHECK: store i32 %{{.*}}, i32 addrspace(1)*
 ; CHECK: store i32 %{{.*}}, i32 addrspace(1)*