Index: clang/lib/CodeGen/TargetInfo.cpp =================================================================== --- clang/lib/CodeGen/TargetInfo.cpp +++ clang/lib/CodeGen/TargetInfo.cpp @@ -449,7 +449,9 @@ // space, an address space conversion may end up as a bitcast. if (auto *C = dyn_cast(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 * Index: clang/test/CodeGenCUDA/builtins-amdgcn.cu =================================================================== --- clang/test/CodeGenCUDA/builtins-amdgcn.cu +++ clang/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: %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; Index: clang/test/CodeGenOpenCL/address-spaces-conversions.cl =================================================================== --- clang/test/CodeGenOpenCL/address-spaces-conversions.cl +++ clang/test/CodeGenOpenCL/address-spaces-conversions.cl @@ -13,7 +13,7 @@ // 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 Index: clang/test/CodeGenOpenCL/atomic-ops-libcall.cl =================================================================== --- clang/test/CodeGenOpenCL/atomic-ops-libcall.cl +++ clang/test/CodeGenOpenCL/atomic-ops-libcall.cl @@ -28,17 +28,17 @@ // 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); @@ -55,28 +55,28 @@ // 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); } Index: clang/test/CodeGenOpenCLCXX/address-space-deduction.cl =================================================================== --- clang/test/CodeGenOpenCLCXX/address-space-deduction.cl +++ clang/test/CodeGenOpenCLCXX/address-space-deduction.cl @@ -29,7 +29,7 @@ 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 Index: clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl =================================================================== --- clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl +++ clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl @@ -94,30 +94,30 @@ // 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* @@ -125,9 +125,9 @@ // 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]] @@ -155,18 +155,18 @@ // 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)* @@ -177,28 +177,28 @@ // 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)* Index: clang/test/CodeGenOpenCLCXX/addrspace-operators.cl =================================================================== --- clang/test/CodeGenOpenCLCXX/addrspace-operators.cl +++ clang/test/CodeGenOpenCLCXX/addrspace-operators.cl @@ -19,11 +19,11 @@ //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; Index: clang/test/CodeGenOpenCLCXX/addrspace-references.cl =================================================================== --- clang/test/CodeGenOpenCLCXX/addrspace-references.cl +++ clang/test/CodeGenOpenCLCXX/addrspace-references.cl @@ -8,7 +8,7 @@ // 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); } Index: clang/test/CodeGenOpenCLCXX/template-address-spaces.cl =================================================================== --- clang/test/CodeGenOpenCLCXX/template-address-spaces.cl +++ clang/test/CodeGenOpenCLCXX/template-address-spaces.cl @@ -13,12 +13,12 @@ // 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 sint;