vzakhari created this revision.
vzakhari added a reviewer: rjmccall.
Herald added subscribers: cfe-commits, jfb, jvesely.
Herald added a project: clang.
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.
Repository:
rC Clang
https://reviews.llvm.org/D63846
Files:
clang/lib/CodeGen/TargetInfo.cpp
clang/test/CodeGenCUDA/builtins-amdgcn.cu
clang/test/CodeGenOpenCL/address-spaces-conversions.cl
clang/test/CodeGenOpenCL/atomic-ops-libcall.cl
clang/test/CodeGenOpenCLCXX/address-space-deduction.cl
clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl
clang/test/CodeGenOpenCLCXX/addrspace-operators.cl
clang/test/CodeGenOpenCLCXX/addrspace-references.cl
clang/test/CodeGenOpenCLCXX/template-address-spaces.cl
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<int> sint;
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/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-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/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/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/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/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/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<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 *
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
