[clang] Preserve names of addrspacecast'ed values.

[vzakhari]

Summary

Attach ".ascast" suffix to a value name when generating addrspacecast for it. This improves IR readability, e.g. for alloca variables, since all users of the variable will be using the addrspacecast value instead of the original alloca.

[vzakhari]

ping

[vzakhari]

Adding more reviewers.

[vzakhari]

ping

[erichkeane]

The change itself doesn't really bother me at all, however the test will fail if we build without preserving names. In order to validate this you'll have to do fno-discard-value-names on a test.

clang/test/CodeGenOpenCL/address-spaces-conversions.cl

16

You probably don't want to remove the wildcard here. If this is built without names being saved, this change will fail.

[rjmccall]

Please don't check IR names in test output. That actually includes anonymous names like %2; these should always be tested with FileCheck variables. I suggest using %.* as the pattern; if you're matching the LHS of an LLVM assignment, that shouldn't have problems with accidentally matching too much.

[vzakhari]

In D63846#1574102, @rjmccall wrote:

Please don't check IR names in test output. That actually includes anonymous names like %2; these should always be tested with FileCheck variables. I suggest using %.* as the pattern; if you're matching the LHS of an LLVM assignment, that shouldn't have problems with accidentally matching too much.

I agree that checking IR names is a bad practice. I can change all these tests so that they either use FileCheck variables or -fno-discard-value-names (e.g. for clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl, which relies on the names for checking) - would that be an appropriate solution?

clang/test/CodeGenOpenCL/address-spaces-conversions.cl

16

In the tests that already use variable names I stick to checking the names, e.g. this test already assumes that the names are preserved (%var_priv in this check), so I am just checking that 'ascast' suffix is added. I guess I will have to add -fno-discard-value-names for this test - do you agree?

[rjmccall]

I don't know what I think about widespread use of -fno-discard-value-names for now; please continue to use FileCheck variables, and we can make a holistic decision about that flag later.

[vzakhari]

In D63846#1574302, @rjmccall wrote:

I don't know what I think about widespread use of -fno-discard-value-names for now; please continue to use FileCheck variables, and we can make a holistic decision about that flag later.

Sorry, I have one particular question about clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl:
Test the address space of 'this' when invoking copy-constructor.
COMMON: [[C1GEN:%c1.ascast[0-9]*]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*

This check seems to rely on %c1 name already. I guess the matching may go off, if we do not use actual names on the right hand side of the assignment. Should I do anything about the right hand side, or just use a generic wildcard on the left hand side?

[rjmccall]

In D63846#1574311, @vzakhari wrote:

In D63846#1574302, @rjmccall wrote:

I don't know what I think about widespread use of -fno-discard-value-names for now; please continue to use FileCheck variables, and we can make a holistic decision about that flag later.

Sorry, I have one particular question about clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl:
Test the address space of 'this' when invoking copy-constructor.
COMMON: [[C1GEN:%c1.ascast[0-9]*]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*

This check seems to rely on %c1 name already. I guess the matching may go off, if we do not use actual names on the right hand side of the assignment. Should I do anything about the right hand side, or just use a generic wildcard on the left hand side?

I forgot that of course you're updating this test already. If you can update the test to use variables without damaging the quality of the test, please do so.

[vzakhari]

I changed the tests to use FileCheck variables whenever possible.

[rjmccall]

Thanks!

clang/lib/CodeGen/TargetInfo.cpp

 
 llvm::Value *TargetCodeGenInfo::performAddrSpaceCast(
     CodeGen::CodeGenFunction &CGF, llvm::Value *Src, LangAS SrcAddr,
     LangAS DestAddr, llvm::Type *DestTy, bool isNonNull) const {
   // Since target may map different address spaces in AST to the same address
   // 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 *
 TargetCodeGenInfo::performAddrSpaceCast(CodeGenModule &CGM, llvm::Constant *Src,
                                         LangAS SrcAddr, LangAS DestAddr,
                                         llvm::Type *DestTy) const {
   // Since target may map different address spaces in AST to the same address
   // space, an address space conversion may end up as a bitcast.

clang/test/CodeGenCUDA/builtins-amdgcn.cu

 // RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm %s -o - | FileCheck %s
 #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: %0 = call i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
+// CHECK: %1 = addrspacecast i8 addrspace(4)* %0 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 %0, i32 0, i32 0, i1 false)
 __global__
 void test_ds_fmax(float src) {
   __shared__ float shared;
   volatile float x = __builtin_amdgcn_ds_fmaxf(&shared, src, 0, 0, false);
 }

clang/test/CodeGenOpenCL/address-spaces-conversions.cl

 // RUN: %clang_cc1 %s -triple x86_64-unknown-linux-gnu -O0 -ffake-address-space-map -cl-std=CL2.0 -emit-llvm -o - | FileCheck %s
 // RUN: %clang_cc1 %s -triple x86_64-unknown-linux-gnu -O0 -cl-std=CL2.0 -emit-llvm -o - | FileCheck --check-prefix=CHECK-NOFAKE %s
 // When -ffake-address-space-map is not used, all addr space mapped to 0 for x86_64.
 
 // test that we generate address space casts everywhere we need conversions of
 // pointers to different address spaces
 
 // CHECK: define void @test
 void test(global int *arg_glob, generic int *arg_gen) {
   int var_priv;
   arg_gen = arg_glob; // implicit cast global -> generic
   // CHECK: %{{[0-9]+}} = addrspacecast i32 addrspace(1)* %{{[0-9]+}} to i32 addrspace(4)*
   // 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: %var_priv.ascast = addrspacecast i32* %var_priv to i32 addrspace(4)*
   // CHECK-NOFAKE-NOT: addrspacecast
 
   arg_glob = (global int *)arg_gen; // explicit cast
   // CHECK: %{{[0-9]+}} = addrspacecast i32 addrspace(4)* %{{[0-9]+}} to i32 addrspace(1)*
   // CHECK-NOFAKE-NOT: addrspacecast
 
   global int *var_glob =
       (global int *)arg_glob; // explicit cast in the same address space

clang/test/CodeGenOpenCL/atomic-ops-libcall.cl

void f(atomic_int *i, global atomic_int *gi, local atomic_int *li, private atomic_int *pi, atomic_uint *ui, int cmp, int order, int scope) {
   // SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_load_4(i8 addrspace(4)* {{%[0-9]+}}, i32 5, i32 1)
   // ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_load_4(i8* {{%[0-9]+}}, i32 5, i32 1)
   x = __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_work_group);
 
   // SPIR: call void @__opencl_atomic_store_4(i8 addrspace(4)* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 1)
   // ARM: call void @__opencl_atomic_store_4(i8* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 1)
   __opencl_atomic_store(i, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: %[[GP:[0-9]+]] = addrspacecast i8 addrspace(1)* {{%[0-9]+}} to i8 addrspace(4)*
+  // SPIR: %[[GP:[.a-z0-9]+]] = addrspacecast i8 addrspace(1)* {{%[0-9]+}} to i8 addrspace(4)*
   // SPIR: call void @__opencl_atomic_store_4(i8 addrspace(4)* %[[GP]], i32 {{%[0-9]+}}, i32 5, i32 1)
   // ARM: call void @__opencl_atomic_store_4(i8* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 1)
   __opencl_atomic_store(gi, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: %[[GP:[0-9]+]] = addrspacecast i8 addrspace(3)* {{%[0-9]+}} to i8 addrspace(4)*
+  // SPIR: %[[GP:[.a-z0-9]+]] = addrspacecast i8 addrspace(3)* {{%[0-9]+}} to i8 addrspace(4)*
   // SPIR: call void @__opencl_atomic_store_4(i8 addrspace(4)* %[[GP]], i32 {{%[0-9]+}}, i32 5, i32 1)
   // ARM: call void @__opencl_atomic_store_4(i8* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 1)
   __opencl_atomic_store(li, 1, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: %[[GP:[0-9]+]] = addrspacecast i8* {{%[0-9]+}} to i8 addrspace(4)*
+  // SPIR: %[[GP:[.a-z0-9]+]] = addrspacecast i8* {{%[0-9]+}} to i8 addrspace(4)*
   // SPIR: call void @__opencl_atomic_store_4(i8 addrspace(4)* %[[GP]], i32 {{%[0-9]+}}, i32 5, i32 1)
   // ARM: call void @__opencl_atomic_store_4(i8* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 1)
   __opencl_atomic_store(pi, 1, memory_order_seq_cst, memory_scope_work_group);
 
   // SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_add_4(i8 addrspace(4)* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 1)
   // ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_add_4(i8* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 1)
   x = __opencl_atomic_fetch_add(i, 3, memory_order_seq_cst, memory_scope_work_group);
 
   // SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_min_4(i8 addrspace(4)* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 1)
   // ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_min_4(i8* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 1)
   x = __opencl_atomic_fetch_min(i, 3, memory_order_seq_cst, memory_scope_work_group);
 
   // SPIR: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_umin_4(i8 addrspace(4)* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 1)
   // ARM: {{%[^ ]*}} = call i32 @__opencl_atomic_fetch_umin_4(i8* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 1)
   x = __opencl_atomic_fetch_min(ui, 3, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* {{%[0-9]+}}, i8 addrspace(4)* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 5, i32 1)
+  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* {{%[0-9]+}}, i8 addrspace(4)* {{%[.a-z0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 5, i32 1)
   // ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8* {{%[0-9]+}}, i8* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 5, i32 1)
   x = __opencl_atomic_compare_exchange_strong(i, &cmp, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* {{%[0-9]+}}, i8 addrspace(4)* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 5, i32 1)
+  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* {{%[0-9]+}}, i8 addrspace(4)* {{%[.a-z0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 5, i32 1)
   // ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8* {{%[0-9]+}}, i8* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 5, i32 1)
   x = __opencl_atomic_compare_exchange_weak(i, &cmp, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
 
-  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* {{%[0-9]+}}, i8 addrspace(4)* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 5, i32 2)
+  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* {{%[0-9]+}}, i8 addrspace(4)* {{%[.a-z0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 5, i32 2)
   // ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8* {{%[0-9]+}}, i8* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 5, i32 2)
   x = __opencl_atomic_compare_exchange_weak(i, &cmp, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_device);
 
-  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* {{%[0-9]+}}, i8 addrspace(4)* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 5, i32 3)
+  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* {{%[0-9]+}}, i8 addrspace(4)* {{%[.a-z0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 5, i32 3)
   // ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8* {{%[0-9]+}}, i8* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 5, i32 3)
   x = __opencl_atomic_compare_exchange_weak(i, &cmp, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_all_svm_devices);
 
 #ifdef cl_khr_subgroups
-  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* {{%[0-9]+}}, i8 addrspace(4)* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 5, i32 4)
+  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* {{%[0-9]+}}, i8 addrspace(4)* {{%[.a-z0-9]+}}, i32 {{%[0-9]+}}, i32 5, i32 5, i32 4)
   x = __opencl_atomic_compare_exchange_weak(i, &cmp, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_sub_group);
 #endif
 
-  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* {{%[0-9]+}}, i8 addrspace(4)* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
+  // SPIR: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8 addrspace(4)* {{%[0-9]+}}, i8 addrspace(4)* {{%[.a-z0-9]+}}, i32 {{%[0-9]+}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
   // ARM: {{%[^ ]*}} = call zeroext i1 @__opencl_atomic_compare_exchange_4(i8* {{%[0-9]+}}, i8* {{%[0-9]+}}, i32 {{%[0-9]+}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
   x = __opencl_atomic_compare_exchange_weak(i, &cmp, 1, order, order, scope);
 }

clang/test/CodeGenOpenCLCXX/address-space-deduction.cl

 //REF: store i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @glob to i32 addrspace(4)*), i32 addrspace(4)* addrspace(1)* @glob_p
 
 //COMMON: define spir_func i32 @_Z3fooi{{P|R}}U3AS4i(i32 %par, i32 addrspace(4)*{{.*}} %par_p)
 int foo(int par, int PTR par_p){
   //COMMON: %loc = alloca i32
   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:%loc.ascast]] = 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
   const __private int PTR loc_p_const = ADR(loc);
 
   // CHECK directives for the following code are located above.
   static int loc_st;
   //REF: store i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @_ZZ3fooiRU3AS4iE6loc_st to i32 addrspace(4)*), i32 addrspace(4)* addrspace(1)* @_ZZ3fooiRU3AS4iE8loc_st_p
   static int PTR loc_st_p = ADR(loc_st);
   extern int loc_ext;
   extern int PTR loc_ext_p;
   (void)loc_ext_p;
   return loc_ext;
 }

clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl

 // COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
 // Test the address space of 'this' when invoking a method using a pointer to the object.
 // COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
 
 // Test the address space of 'this' when invoking a method that is declared in the file contex.
 // 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:%c1.ascast[0-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:%c2.ascast[0-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:%c1.ascast[0-9]*]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// COMMON:  [[C2GEN:%c2.ascast[0-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:%c1.ascast[0-9]*]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// COMMON: [[C2GEN:%c2.ascast[0-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:%c4.ascast[0-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*
 // IMPL: [[CALLVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CALL]] to i8 addrspace(4)*
 // 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:%c5.ascast[0-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]]
 
 // Tests address space of inline members
 //COMMON: @_ZNU3AS41C3getEv(%class.C addrspace(4)* %this)
 //COMMON: @_ZNU3AS41CplERU3AS4KS_(%class.C* noalias sret %agg.result, %class.C addrspace(4)* %this
 #define TEST(AS)             \

 
 // Test that we don't initialize an object in local address space.
 // EXPL-NOT: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localE1c to %class.C addrspace(4)*))
 
 // Test the address space of 'this' when invoking a method.
 // 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:%c1.ascast[0-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:%c2.ascast[0-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:%c1.ascast[0-9]*]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// COMMON: [[C2GEN:%c2.ascast[0-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(__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:%c.ascast[0-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:%c.ascast[0-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:%c1.ascast[0-9]*]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// COMMON: [[CGEN:%c.ascast[0-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:%c2.ascast[0-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:%c1.ascast[0-9]*]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// COMMON: [[C2GEN:%c2.ascast[0-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]]

clang/test/CodeGenOpenCLCXX/addrspace-operators.cl

 };
 
 __global E globE;
 volatile __global int globVI;
 __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:%c.ascast[0-9]*]] = addrspacecast %class.C* %c 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:%c.ascast[0-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;
   //CHECK: store i32 1, i32* %e
   e = b;
   //CHECK: store i32 0, i32 addrspace(1)* @globE
   globE = a;
   //CHECK: store i32 %or, i32 addrspace(1)* @globI

clang/test/CodeGenOpenCLCXX/addrspace-references.cl

 //RUN: %clang_cc1 %s -cl-std=c++ -triple spir -emit-llvm -o - | FileCheck %s
 
 int bar(const unsigned int &i);
 // CHECK-LABEL: define spir_func void @_Z3foov() 
 void foo() {
   // The generic addr space reference parameter object will be bound
   // to a temporary value allocated in private addr space. We need an
   // 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);
 }

clang/test/CodeGenOpenCLCXX/template-address-spaces.cl

 // RUN: %clang_cc1 -cl-std=c++ %s -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck %s
 
 template <typename T>
 struct S{
   T a;
   T foo();
 };
 
 template<typename T>
 T S<T>::foo() { return a;}
 
 // CHECK: %struct.S = type { i32 }
 // 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:%sint.ascast[0-9]*]] = addrspacecast %struct.S* %sint to %struct.S addrspace(4)*
+// CHECK:  %call = call i32 @_ZNU3AS41SIiE3fooEv(%struct.S addrspace(4)* [[A1]]) #1
+// CHECK:  [[A2:%sintptr.ascast[0-9]*]] = addrspacecast %struct.S.0* %sintptr to %struct.S.0 addrspace(4)*
+// CHECK:  %call1 = call i32 addrspace(4)* @_ZNU3AS41SIPU3AS4iE3fooEv(%struct.S.0 addrspace(4)* [[A2]]) #1
+// CHECK:  [[A3:%sintptrgl.ascast[0-9]*]] = addrspacecast %struct.S.1* %sintptrgl 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;
   S<int*> sintptr;
   S<__global int*> sintptrgl;
 
   sint.foo();
   sintptr.foo();
   sintptrgl.foo();
 }