[clang] Preserve names of addrspacecast'ed values.

Differential Revision: https://reviews.llvm.org/D63846

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@365666 91177308-0d34-0410-b5e6-96231b3b80d8

lib/CodeGen/TargetInfo.cpp

@@ -449,7 +449,9 @@ llvm::Value *TargetCodeGenInfo::performAddrSpaceCast(
   // space, an address space conversion may end up as a bitcast.
   if (auto *C = dyn_cast<llvm::Constant>(Src))
     return performAddrSpaceCast(CGF.CGM, C, SrcAddr, DestAddr, DestTy);
-  return CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(Src, DestTy);
+  // Try to preserve the source's name to make IR more readable.
+  return CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
+      Src, DestTy, Src->hasName() ? Src->getName() + ".ascast" : "");
 }
 
 llvm::Constant *

test/CodeGenCUDA/builtins-amdgcn.cu

@@ -2,15 +2,15 @@
 #include "Inputs/cuda.h"
 
 // CHECK-LABEL: @_Z16use_dispatch_ptrPi(
-// CHECK: %2 = call i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
-// CHECK: %3 = addrspacecast i8 addrspace(4)* %2 to i8 addrspace(4)**
+// CHECK: %[[PTR:.*]] = call i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
+// CHECK: %{{.*}} = addrspacecast i8 addrspace(4)* %[[PTR]] to i8 addrspace(4)**
 __global__ void use_dispatch_ptr(int* out) {
   const int* dispatch_ptr = (const int*)__builtin_amdgcn_dispatch_ptr();
   *out = *dispatch_ptr;
 }
 
 // CHECK-LABEL: @_Z12test_ds_fmaxf(
-// CHECK: call float @llvm.amdgcn.ds.fmax(float addrspace(3)* @_ZZ12test_ds_fmaxfE6shared, float %2, i32 0, i32 0, i1 false)
+// CHECK: call float @llvm.amdgcn.ds.fmax(float addrspace(3)* @_ZZ12test_ds_fmaxfE6shared, float %{{[^,]*}}, i32 0, i32 0, i1 false)
 __global__
 void test_ds_fmax(float src) {
   __shared__ float shared;

test/CodeGenOpenCL/address-spaces-conversions.cl

@@ -13,7 +13,7 @@ void test(global int *arg_glob, generic int *arg_gen) {
   // CHECK-NOFAKE-NOT: addrspacecast
 
   arg_gen = &var_priv; // implicit cast with obtaining adr, private -> generic
-  // CHECK: %{{[0-9]+}} = addrspacecast i32* %var_priv to i32 addrspace(4)*
+  // CHECK: %{{[._a-z0-9]+}} = addrspacecast i32* %{{[._a-z0-9]+}} to i32 addrspace(4)*
   // CHECK-NOFAKE-NOT: addrspacecast
 
   arg_glob = (global int *)arg_gen; // explicit cast

test/CodeGenOpenCLCXX/address-space-deduction.cl

@@ -29,7 +29,7 @@ int foo(int par, int PTR par_p){
   int loc;
   //COMMON: %loc_p = alloca i32 addrspace(4)*
   //COMMON: %loc_p_const = alloca i32*
-  //COMMON: [[GAS:%[0-9]+]] = addrspacecast i32* %loc to i32 addrspace(4)*
+  //COMMON: [[GAS:%[._a-z0-9]*]] = addrspacecast i32* %loc to i32 addrspace(4)*
   //COMMON: store i32 addrspace(4)* [[GAS]], i32 addrspace(4)** %loc_p
   int PTR loc_p = ADR(loc);
   //COMMON: store i32* %loc, i32** %loc_p_const

test/CodeGenOpenCLCXX/addrspace-of-this.cl

@@ -94,30 +94,30 @@ __kernel void test__global() {
 // COMMON: call i32 @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
 
 // Test the address space of 'this' when invoking copy-constructor.
-// COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
 // IMPL: [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8*
 // IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C1VOID]], i8 addrspace(4)* {{.*}}addrspacecast (i8 addrspace(1)* bitcast (%class.C addrspace(1)* @c to i8 addrspace(1)*) to i8 addrspace(4)*)
 // EXPL: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
 
 // Test the address space of 'this' when invoking a constructor.
-// EXPL:   [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// EXPL:   [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
 // EXPL:   call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
 
 // Test the address space of 'this' when invoking assignment operator.
-// COMMON:  [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// COMMON:  [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// COMMON:  [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// COMMON:  [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
 // EXPL: call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]])
 // IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)*
 // IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)*
 // IMPL: call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* {{.*}}[[C2GENVOID]], i8 addrspace(4)* {{.*}}[[C1GENVOID]]
 
 // Test the address space of 'this' when invoking the operator+
-// COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// COMMON: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
 // COMMON: call void @_ZNU3AS41CplERU3AS4KS_(%class.C* sret %c3, %class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[C2GEN]])
 
 // Test the address space of 'this' when invoking the move constructor
-// COMMON: [[C4GEN:%[0-9]+]] = addrspacecast %class.C* %c4 to %class.C addrspace(4)*
+// COMMON: [[C4GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c4 to %class.C addrspace(4)*
 // COMMON: [[CALL:%call[0-9]+]] = call spir_func dereferenceable(4) %class.C addrspace(4)* @_Z3foov()
 // EXPL: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C4GEN]], %class.C addrspace(4)* dereferenceable(4) [[CALL]])
 // IMPL: [[C4VOID:%[0-9]+]] = bitcast %class.C* %c4 to i8*
@@ -125,9 +125,9 @@ __kernel void test__global() {
 // IMPL:  call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C4VOID]], i8 addrspace(4)* {{.*}}[[CALLVOID]]
 
 // Test the address space of 'this' when invoking the move assignment
-// COMMON: [[C5GEN:%[0-9]+]] = addrspacecast %class.C* %c5 to %class.C addrspace(4)*
+// COMMON: [[C5GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c5 to %class.C addrspace(4)*
 // COMMON: [[CALL:%call[0-9]+]] = call spir_func dereferenceable(4) %class.C addrspace(4)* @_Z3foov()
-// EXPL: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C5GEN:%[0-9]+]], %class.C addrspace(4)* dereferenceable(4) [[CALL]])
+// EXPL: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C5GEN]], %class.C addrspace(4)* dereferenceable(4) [[CALL]])
 // IMPL: [[C5VOID:%[0-9]+]] = bitcast %class.C* %c5 to i8*
 // IMPL: [[CALLVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CALL]] to i8 addrspace(4)*
 // IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C5VOID]], i8 addrspace(4)* {{.*}}[[CALLVOID]]
@@ -155,18 +155,18 @@ TEST(__local)
 // COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localE1c to %class.C addrspace(4)*))
 
 // Test the address space of 'this' when invoking copy-constructor.
-// COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
 // EXPL: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(3)* @_ZZ11test__localE1c to %class.C addrspace(4)*))
 // IMPL:  [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8*
 // IMPL:  call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C1VOID]], i8 addrspace(4)* {{.*}}addrspacecast (i8 addrspace(3)* bitcast (%class.C addrspace(3)* @_ZZ11test__localE1c to i8 addrspace(3)*) to i8 addrspace(4)*), i32 4, i1 false)
 
 // Test the address space of 'this' when invoking a constructor.
-// EXPL: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// EXPL: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
 // EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
 
 // Test the address space of 'this' when invoking assignment operator.
-// COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// COMMON: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
 // EXPL: call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]])
 // IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)*
 // IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)*
@@ -177,28 +177,28 @@ TEST(__private)
 // CHECK-LABEL: @test__private
 
 // Test the address space of 'this' when invoking a constructor for an object in non-default address space
-// EXPL: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// EXPL: [[CGEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
 // EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[CGEN]])
 
 // Test the address space of 'this' when invoking a method.
-// COMMON: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// COMMON: [[CGEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
 // COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]])
 
 // Test the address space of 'this' when invoking a copy-constructor.
-// COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// COMMON: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// COMMON: [[CGEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
 // EXPL: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[CGEN]])
 // IMPL: [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8*
 // IMPL: [[CGENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CGEN]] to i8 addrspace(4)*
 // IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C1VOID]], i8 addrspace(4)* {{.*}}[[CGENVOID]]
 
 // Test the address space of 'this' when invoking a constructor.
-// EXPL: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// EXPL: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
 // EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
 
 // Test the address space of 'this' when invoking a copy-assignment.
-// COMMON: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// COMMON: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
 // EXPL: call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]])
 // IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)*
 // IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)*

test/CodeGenOpenCLCXX/addrspace-operators.cl

@@ -19,11 +19,11 @@ __global int globI;
 //CHECK-LABEL: define spir_func void @_Z3barv()
 void bar() {
   C c;
-  //CHECK: addrspacecast %class.C* %c to %class.C addrspace(4)*
-  //CHECK: call void @_ZNU3AS41C6AssignE1E(%class.C addrspace(4)* %{{[0-9]+}}, i32 0)
+  //CHECK: [[A1:%[.a-z0-9]+]] = addrspacecast %class.C* [[C:%[a-z0-9]+]] to %class.C addrspace(4)*
+  //CHECK: call void @_ZNU3AS41C6AssignE1E(%class.C addrspace(4)* [[A1]], i32 0)
   c.Assign(a);
-  //CHECK: addrspacecast %class.C* %c to %class.C addrspace(4)*
-  //CHECK: call void @_ZNU3AS41C8OrAssignE1E(%class.C addrspace(4)* %{{[0-9]+}}, i32 0)
+  //CHECK: [[A2:%[.a-z0-9]+]] = addrspacecast %class.C* [[C]] to %class.C addrspace(4)*
+  //CHECK: call void @_ZNU3AS41C8OrAssignE1E(%class.C addrspace(4)* [[A2]], i32 0)
   c.OrAssign(a);
 
   E e;
@@ -35,19 +35,33 @@ void bar() {
   globI |= b;
   //CHECK: store i32 %add, i32 addrspace(1)* @globI
   globI += a;
-  //CHECK: store volatile i32 %and, i32 addrspace(1)* @globVI
+  //CHECK: [[GVIV1:%[0-9]+]] = load volatile i32, i32 addrspace(1)* @globVI
+  //CHECK: [[AND:%[a-z0-9]+]] = and i32 [[GVIV1]], 1
+  //CHECK: store volatile i32 [[AND]], i32 addrspace(1)* @globVI
   globVI &= b;
-  //CHECK: store volatile i32 %sub, i32 addrspace(1)* @globVI
+  //CHECK: [[GVIV2:%[0-9]+]] = load volatile i32, i32 addrspace(1)* @globVI
+  //CHECK: [[SUB:%[a-z0-9]+]] = sub nsw i32 [[GVIV2]], 0
+  //CHECK: store volatile i32 [[SUB]], i32 addrspace(1)* @globVI
   globVI -= a;
 }
 
-//CHECK: define linkonce_odr void @_ZNU3AS41C6AssignE1E(%class.C addrspace(4)* %this, i32 %e)
-//CHECK: [[E:%[0-9]+]] = load i32, i32* %e.addr
-//CHECK: %me = getelementptr inbounds %class.C, %class.C addrspace(4)* %this1, i32 0, i32 0
-//CHECK: store i32 [[E]], i32 addrspace(4)* %me
+//CHECK: define linkonce_odr void @_ZNU3AS41C6AssignE1E(%class.C addrspace(4)*{{[ %a-z0-9]*}}, i32{{[ %a-z0-9]*}})
+//CHECK: [[THIS_ADDR:%[.a-z0-9]+]] = alloca %class.C addrspace(4)
+//CHECK: [[E_ADDR:%[.a-z0-9]+]] = alloca i32
+//CHECK: store %class.C addrspace(4)* {{%[a-z0-9]+}}, %class.C addrspace(4)** [[THIS_ADDR]]
+//CHECK: store i32 {{%[a-z0-9]+}}, i32* [[E_ADDR]]
+//CHECK: [[THIS1:%[.a-z0-9]+]] = load %class.C addrspace(4)*, %class.C addrspace(4)** [[THIS_ADDR]]
+//CHECK: [[E:%[0-9]+]] = load i32, i32* [[E_ADDR]]
+//CHECK: [[ME:%[a-z0-9]+]] = getelementptr inbounds %class.C, %class.C addrspace(4)* [[THIS1]], i32 0, i32 0
+//CHECK: store i32 [[E]], i32 addrspace(4)* [[ME]]
 
-//CHECK define linkonce_odr void @_ZNU3AS41C8OrAssignE1E(%class.C addrspace(4)* %this, i32 %e)
-//CHECK: [[E:%[0-9]+]] = load i32, i32* %e.addr
-//CHECK: %mi = getelementptr inbounds %class.C, %class.C addrspace(4)* %this1, i32 0, i32 1
-//CHECK: [[MI:%[0-9]+]] = load i32, i32 addrspace(4)* %mi
+//CHECK: define linkonce_odr void @_ZNU3AS41C8OrAssignE1E(%class.C addrspace(4)*{{[ %a-z0-9]*}}, i32{{[ %a-z0-9]*}})
+//CHECK: [[THIS_ADDR:%[.a-z0-9]+]] = alloca %class.C addrspace(4)
+//CHECK: [[E_ADDR:%[.a-z0-9]+]] = alloca i32
+//CHECK: store %class.C addrspace(4)* {{%[a-z0-9]+}}, %class.C addrspace(4)** [[THIS_ADDR]]
+//CHECK: store i32 {{%[a-z0-9]+}}, i32* [[E_ADDR]]
+//CHECK: [[THIS1:%[.a-z0-9]+]] = load %class.C addrspace(4)*, %class.C addrspace(4)** [[THIS_ADDR]]
+//CHECK: [[E:%[0-9]+]] = load i32, i32* [[E_ADDR]]
+//CHECK: [[MI_GEP:%[a-z0-9]+]] = getelementptr inbounds %class.C, %class.C addrspace(4)* [[THIS1]], i32 0, i32 1
+//CHECK: [[MI:%[0-9]+]] = load i32, i32 addrspace(4)* [[MI_GEP]]
 //CHECK: %or = or i32 [[MI]], [[E]]

test/CodeGenOpenCLCXX/addrspace-references.cl

@@ -8,7 +8,7 @@ void foo() {
   // addrspacecast before passing the value to the function.
   // CHECK: [[REF:%.*]] = alloca i32
   // CHECK: store i32 1, i32* [[REF]]
-  // CHECK: [[REG:%[0-9]+]] = addrspacecast i32* [[REF]] to i32 addrspace(4)*
+  // CHECK: [[REG:%[.a-z0-9]+]] = addrspacecast i32* [[REF]] to i32 addrspace(4)*
   // CHECK: call spir_func i32 @_Z3barRU3AS4Kj(i32 addrspace(4)* dereferenceable(4) [[REG]])
   bar(1);
 }

test/CodeGenOpenCLCXX/template-address-spaces.cl

@@ -13,12 +13,12 @@ T S<T>::foo() { return a;}
 // CHECK: %struct.S.0 = type { i32 addrspace(4)* }
 // CHECK: %struct.S.1 = type { i32 addrspace(1)* }
 
-// CHECK:  %0 = addrspacecast %struct.S* %sint to %struct.S addrspace(4)*
-// CHECK:  %call = call i32 @_ZNU3AS41SIiE3fooEv(%struct.S addrspace(4)* %0) #1
-// CHECK:  %1 = addrspacecast %struct.S.0* %sintptr to %struct.S.0 addrspace(4)*
-// CHECK:  %call1 = call i32 addrspace(4)* @_ZNU3AS41SIPU3AS4iE3fooEv(%struct.S.0 addrspace(4)* %1) #1
-// CHECK:  %2 = addrspacecast %struct.S.1* %sintptrgl to %struct.S.1 addrspace(4)*
-// CHECK:  %call2 = call i32 addrspace(1)* @_ZNU3AS41SIPU3AS1iE3fooEv(%struct.S.1 addrspace(4)* %2) #1
+// CHECK:  [[A1:%[.a-z0-9]+]] = addrspacecast %struct.S* %{{[a-z0-9]+}} to %struct.S addrspace(4)*
+// CHECK:  %call = call i32 @_ZNU3AS41SIiE3fooEv(%struct.S addrspace(4)* [[A1]]) #1
+// CHECK:  [[A2:%[.a-z0-9]+]] = addrspacecast %struct.S.0* %{{[a-z0-9]+}} to %struct.S.0 addrspace(4)*
+// CHECK:  %call1 = call i32 addrspace(4)* @_ZNU3AS41SIPU3AS4iE3fooEv(%struct.S.0 addrspace(4)* [[A2]]) #1
+// CHECK:  [[A3:%[.a-z0-9]+]] = addrspacecast %struct.S.1* %{{[a-z0-9]+}} to %struct.S.1 addrspace(4)*
+// CHECK:  %call2 = call i32 addrspace(1)* @_ZNU3AS41SIPU3AS1iE3fooEv(%struct.S.1 addrspace(4)* [[A3]]) #1
 
 void bar(){
   S<int> sint;