r367103 - Make the CXXABIs respect the target's default calling convention.
Erich Keane via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 26 05:36:12 PDT 2019
Author: erichkeane
Date: Fri Jul 26 05:36:12 2019
New Revision: 367103
URL: http://llvm.org/viewvc/llvm-project?rev=367103&view=rev
Log:
Make the CXXABIs respect the target's default calling convention.
SPIR targets need to have all functions be SPIR calling convention,
however the CXXABIs were just returning CC_C in all non-'this-CC' cases.
https://reviews.llvm.org/D65294
Modified:
cfe/trunk/lib/AST/ItaniumCXXABI.cpp
cfe/trunk/lib/AST/MicrosoftCXXABI.cpp
cfe/trunk/test/CodeGenOpenCLCXX/addrspace-derived-base.cl
cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl
cfe/trunk/test/CodeGenOpenCLCXX/addrspace-operators.cl
cfe/trunk/test/CodeGenOpenCLCXX/addrspace-with-class.cl
cfe/trunk/test/CodeGenOpenCLCXX/method-overload-address-space.cl
cfe/trunk/test/CodeGenOpenCLCXX/template-address-spaces.cl
Modified: cfe/trunk/lib/AST/ItaniumCXXABI.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ItaniumCXXABI.cpp?rev=367103&r1=367102&r2=367103&view=diff
==============================================================================
--- cfe/trunk/lib/AST/ItaniumCXXABI.cpp (original)
+++ cfe/trunk/lib/AST/ItaniumCXXABI.cpp Fri Jul 26 05:36:12 2019
@@ -177,7 +177,7 @@ public:
if (!isVariadic && T.isWindowsGNUEnvironment() &&
T.getArch() == llvm::Triple::x86)
return CC_X86ThisCall;
- return CC_C;
+ return Context.getTargetInfo().getDefaultCallingConv();
}
// We cheat and just check that the class has a vtable pointer, and that it's
Modified: cfe/trunk/lib/AST/MicrosoftCXXABI.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/MicrosoftCXXABI.cpp?rev=367103&r1=367102&r2=367103&view=diff
==============================================================================
--- cfe/trunk/lib/AST/MicrosoftCXXABI.cpp (original)
+++ cfe/trunk/lib/AST/MicrosoftCXXABI.cpp Fri Jul 26 05:36:12 2019
@@ -82,7 +82,7 @@ public:
if (!isVariadic &&
Context.getTargetInfo().getTriple().getArch() == llvm::Triple::x86)
return CC_X86ThisCall;
- return CC_C;
+ return Context.getTargetInfo().getDefaultCallingConv();
}
bool isNearlyEmpty(const CXXRecordDecl *RD) const override {
Modified: cfe/trunk/test/CodeGenOpenCLCXX/addrspace-derived-base.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCLCXX/addrspace-derived-base.cl?rev=367103&r1=367102&r2=367103&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCLCXX/addrspace-derived-base.cl (original)
+++ cfe/trunk/test/CodeGenOpenCLCXX/addrspace-derived-base.cl Fri Jul 26 05:36:12 2019
@@ -12,11 +12,11 @@ public:
void foo() {
D d;
//CHECK: addrspacecast %class.D* %d to %class.D addrspace(4)*
- //CHECK: call i32 @_ZNU3AS41D5getmbEv(%class.D addrspace(4)*
+ //CHECK: call spir_func i32 @_ZNU3AS41D5getmbEv(%class.D addrspace(4)*
d.getmb();
}
//Derived and Base are in the same address space.
-//CHECK: define linkonce_odr i32 @_ZNU3AS41D5getmbEv(%class.D addrspace(4)* %this)
+//CHECK: define linkonce_odr spir_func i32 @_ZNU3AS41D5getmbEv(%class.D addrspace(4)* %this)
//CHECK: bitcast %class.D addrspace(4)* %this1 to %struct.B addrspace(4)*
Modified: cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl?rev=367103&r1=367102&r2=367103&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl (original)
+++ cfe/trunk/test/CodeGenOpenCLCXX/addrspace-of-this.cl Fri Jul 26 05:36:12 2019
@@ -81,32 +81,32 @@ __kernel void test__global() {
// COMMON: @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* %this)
// EXPL-LABEL: @__cxx_global_var_init()
-// EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+// EXPL: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
// COMMON-LABEL: @test__global()
// Test the address space of 'this' when invoking a method.
-// COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+// COMMON: call spir_func 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)*))
+// COMMON: call spir_func 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)*))
+// COMMON: call spir_func 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:%[.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)*))
+// EXPL: call spir_func 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:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
+// EXPL: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
// Test the address space of 'this' when invoking assignment operator.
// 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]])
+// EXPL: call spir_func 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]]
@@ -114,12 +114,12 @@ __kernel void test__global() {
// Test the address space of 'this' when invoking the operator+
// 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]])
+// COMMON: call spir_func 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:%[.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]])
+// EXPL: call spir_func 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]]
@@ -127,7 +127,7 @@ __kernel void test__global() {
// Test the address space of 'this' when invoking the move assignment
// 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]], %class.C addrspace(4)* dereferenceable(4) [[CALL]])
+// EXPL: call spir_func 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]]
@@ -149,25 +149,25 @@ TEST(__local)
// COMMON-LABEL: @test__local
// 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)*))
+// EXPL-NOT: call spir_func 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)*))
+// COMMON: call spir_func 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:%[.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)*))
+// EXPL: call spir_func 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:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
+// EXPL: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
// Test the address space of 'this' when invoking assignment operator.
// 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]])
+// EXPL: call spir_func 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]]
@@ -178,28 +178,28 @@ TEST(__private)
// Test the address space of 'this' when invoking a constructor for an object in non-default address space
// EXPL: [[CGEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
-// EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[CGEN]])
+// EXPL: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[CGEN]])
// Test the address space of 'this' when invoking a method.
// COMMON: [[CGEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
-// COMMON: call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]])
+// COMMON: call spir_func i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]])
// Test the address space of 'this' when invoking a copy-constructor.
// 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]])
+// EXPL: call spir_func 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:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// EXPL: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
+// EXPL: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
// Test the address space of 'this' when invoking a copy-assignment.
// 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]])
+// EXPL: call spir_func 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]]
Modified: cfe/trunk/test/CodeGenOpenCLCXX/addrspace-operators.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCLCXX/addrspace-operators.cl?rev=367103&r1=367102&r2=367103&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCLCXX/addrspace-operators.cl (original)
+++ cfe/trunk/test/CodeGenOpenCLCXX/addrspace-operators.cl Fri Jul 26 05:36:12 2019
@@ -20,10 +20,10 @@ __global int globI;
void bar() {
C c;
//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)
+ //CHECK: call spir_func void @_ZNU3AS41C6AssignE1E(%class.C addrspace(4)* [[A1]], i32 0)
c.Assign(a);
//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)
+ //CHECK: call spir_func void @_ZNU3AS41C8OrAssignE1E(%class.C addrspace(4)* [[A2]], i32 0)
c.OrAssign(a);
E e;
@@ -45,7 +45,7 @@ void bar() {
globVI -= a;
}
-//CHECK: define linkonce_odr void @_ZNU3AS41C6AssignE1E(%class.C addrspace(4)*{{[ %a-z0-9]*}}, i32{{[ %a-z0-9]*}})
+//CHECK: define linkonce_odr spir_func 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]]
@@ -55,7 +55,7 @@ void bar() {
//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)*{{[ %a-z0-9]*}}, i32{{[ %a-z0-9]*}})
+//CHECK: define linkonce_odr spir_func 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]]
Modified: cfe/trunk/test/CodeGenOpenCLCXX/addrspace-with-class.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCLCXX/addrspace-with-class.cl?rev=367103&r1=367102&r2=367103&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCLCXX/addrspace-with-class.cl (original)
+++ cfe/trunk/test/CodeGenOpenCLCXX/addrspace-with-class.cl Fri Jul 26 05:36:12 2019
@@ -23,17 +23,17 @@ __constant MyType const2(2);
// CHECK: @glob = addrspace(1) global %struct.MyType zeroinitializer
MyType glob(1);
-// CHECK: call void @_ZNU3AS26MyTypeC1Ei(%struct.MyType addrspace(2)* @const1, i32 1)
-// CHECK: call void @_ZNU3AS26MyTypeC1Ei(%struct.MyType addrspace(2)* @const2, i32 2)
-// CHECK: call void @_ZNU3AS46MyTypeC1Ei(%struct.MyType addrspace(4)* addrspacecast (%struct.MyType addrspace(1)* @glob to %struct.MyType addrspace(4)*), i32 1)
+// CHECK: call spir_func void @_ZNU3AS26MyTypeC1Ei(%struct.MyType addrspace(2)* @const1, i32 1)
+// CHECK: call spir_func void @_ZNU3AS26MyTypeC1Ei(%struct.MyType addrspace(2)* @const2, i32 2)
+// CHECK: call spir_func void @_ZNU3AS46MyTypeC1Ei(%struct.MyType addrspace(4)* addrspacecast (%struct.MyType addrspace(1)* @glob to %struct.MyType addrspace(4)*), i32 1)
// CHECK-LABEL: define spir_kernel void @fooGlobal()
kernel void fooGlobal() {
- // CHECK: call i32 @_ZNU3AS46MyType3barEv(%struct.MyType addrspace(4)* addrspacecast (%struct.MyType addrspace(1)* @glob to %struct.MyType addrspace(4)*))
+ // CHECK: call spir_func i32 @_ZNU3AS46MyType3barEv(%struct.MyType addrspace(4)* addrspacecast (%struct.MyType addrspace(1)* @glob to %struct.MyType addrspace(4)*))
glob.bar();
- // CHECK: call i32 @_ZNU3AS26MyType3barEv(%struct.MyType addrspace(2)* @const1)
+ // CHECK: call spir_func i32 @_ZNU3AS26MyType3barEv(%struct.MyType addrspace(2)* @const1)
const1.bar();
- // CHECK: call void @_ZNU3AS26MyTypeD1Ev(%struct.MyType addrspace(2)* @const1)
+ // CHECK: call spir_func void @_ZNU3AS26MyTypeD1Ev(%struct.MyType addrspace(2)* @const1)
const1.~MyType();
}
@@ -41,19 +41,19 @@ kernel void fooGlobal() {
kernel void fooLocal() {
// CHECK: [[VAR:%.*]] = alloca %struct.MyType
// CHECK: [[REG:%.*]] = addrspacecast %struct.MyType* [[VAR]] to %struct.MyType addrspace(4)*
- // CHECK: call void @_ZNU3AS46MyTypeC1Ei(%struct.MyType addrspace(4)* [[REG]], i32 3)
+ // CHECK: call spir_func void @_ZNU3AS46MyTypeC1Ei(%struct.MyType addrspace(4)* [[REG]], i32 3)
MyType myLocal(3);
// CHECK: [[REG:%.*]] = addrspacecast %struct.MyType* [[VAR]] to %struct.MyType addrspace(4)*
- // CHECK: call i32 @_ZNU3AS46MyType3barEv(%struct.MyType addrspace(4)* [[REG]])
+ // CHECK: call spir_func i32 @_ZNU3AS46MyType3barEv(%struct.MyType addrspace(4)* [[REG]])
myLocal.bar();
// CHECK: [[REG:%.*]] = addrspacecast %struct.MyType* [[VAR]] to %struct.MyType addrspace(4)*
- // CHECK: call void @_ZNU3AS46MyTypeD1Ev(%struct.MyType addrspace(4)* [[REG]])
+ // CHECK: call spir_func void @_ZNU3AS46MyTypeD1Ev(%struct.MyType addrspace(4)* [[REG]])
}
// Ensure all members are defined for all the required address spaces.
-// CHECK-DEFINITIONS-DAG: define linkonce_odr void @_ZNU3AS26MyTypeC1Ei(%struct.MyType addrspace(2)* %this, i32 %i)
-// CHECK-DEFINITIONS-DAG: define linkonce_odr void @_ZNU3AS46MyTypeC1Ei(%struct.MyType addrspace(4)* %this, i32 %i)
-// CHECK-DEFINITIONS-DAG: define linkonce_odr void @_ZNU3AS26MyTypeD1Ev(%struct.MyType addrspace(2)* %this)
-// CHECK-DEFINITIONS-DAG: define linkonce_odr void @_ZNU3AS46MyTypeD1Ev(%struct.MyType addrspace(4)* %this)
-// CHECK-DEFINITIONS-DAG: define linkonce_odr i32 @_ZNU3AS26MyType3barEv(%struct.MyType addrspace(2)* %this)
-// CHECK-DEFINITIONS-DAG: define linkonce_odr i32 @_ZNU3AS46MyType3barEv(%struct.MyType addrspace(4)* %this)
+// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func void @_ZNU3AS26MyTypeC1Ei(%struct.MyType addrspace(2)* %this, i32 %i)
+// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func void @_ZNU3AS46MyTypeC1Ei(%struct.MyType addrspace(4)* %this, i32 %i)
+// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func void @_ZNU3AS26MyTypeD1Ev(%struct.MyType addrspace(2)* %this)
+// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func void @_ZNU3AS46MyTypeD1Ev(%struct.MyType addrspace(4)* %this)
+// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func i32 @_ZNU3AS26MyType3barEv(%struct.MyType addrspace(2)* %this)
+// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func i32 @_ZNU3AS46MyType3barEv(%struct.MyType addrspace(4)* %this)
Modified: cfe/trunk/test/CodeGenOpenCLCXX/method-overload-address-space.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCLCXX/method-overload-address-space.cl?rev=367103&r1=367102&r2=367103&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCLCXX/method-overload-address-space.cl (original)
+++ cfe/trunk/test/CodeGenOpenCLCXX/method-overload-address-space.cl Fri Jul 26 05:36:12 2019
@@ -15,21 +15,21 @@ __kernel void k() {
__global C &c_ref = c1;
__global C *c_ptr;
- // CHECK: call void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)*
+ // CHECK: call spir_func void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)*
c1.foo();
- // CHECK: call void @_ZNU3AS31C3fooEv(%struct.C addrspace(3)*
+ // CHECK: call spir_func void @_ZNU3AS31C3fooEv(%struct.C addrspace(3)*
c2.foo();
- // CHECK: call void @_ZNU3AS41C3fooEv(%struct.C addrspace(4)*
+ // CHECK: call spir_func void @_ZNU3AS41C3fooEv(%struct.C addrspace(4)*
c3.foo();
- // CHECK: call void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)*
+ // CHECK: call spir_func void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)*
c_ptr->foo();
- // CHECK: void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)*
+ // CHECK: spir_func void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)*
c_ref.foo();
- // CHECK: call void @_ZNU3AS41C3barEv(%struct.C addrspace(4)* addrspacecast (%struct.C addrspace(1)* @c1 to %struct.C addrspace(4)*))
+ // CHECK: call spir_func void @_ZNU3AS41C3barEv(%struct.C addrspace(4)* addrspacecast (%struct.C addrspace(1)* @c1 to %struct.C addrspace(4)*))
c1.bar();
//FIXME: Doesn't compile yet
//c_ptr->bar();
- // CHECK: call void @_ZNU3AS41C3barEv(%struct.C addrspace(4)* addrspacecast (%struct.C addrspace(1)* @c1 to %struct.C addrspace(4)*))
+ // CHECK: call spir_func void @_ZNU3AS41C3barEv(%struct.C addrspace(4)* addrspacecast (%struct.C addrspace(1)* @c1 to %struct.C addrspace(4)*))
c_ref.bar();
}
Modified: cfe/trunk/test/CodeGenOpenCLCXX/template-address-spaces.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCLCXX/template-address-spaces.cl?rev=367103&r1=367102&r2=367103&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCLCXX/template-address-spaces.cl (original)
+++ cfe/trunk/test/CodeGenOpenCLCXX/template-address-spaces.cl Fri Jul 26 05:36:12 2019
@@ -14,11 +14,11 @@ T S<T>::foo() { return a;}
// CHECK: %struct.S.1 = type { i32 addrspace(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: %call = call spir_func 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: %call1 = call spir_func 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
+// CHECK: %call2 = call spir_func i32 addrspace(1)* @_ZNU3AS41SIPU3AS1iE3fooEv(%struct.S.1 addrspace(4)* [[A3]]) #1
void bar(){
S<int> sint;
More information about the cfe-commits
mailing list