Make the CXXABIs respect the target's default calling convention.

[erichkeane]

Summary

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.

[Anastasia]

LGTM! Thanks!

This should probably go to release 9.0?

[erichkeane]

Committed, 9.0 release request here: https://bugs.llvm.org/show_bug.cgi?id=42774

clang/lib/AST/ItaniumCXXABI.cpp

getMemberPointerInfo(const MemberPointerType *MPT) const override {
     return MPI;
   }
 
   CallingConv getDefaultMethodCallConv(bool isVariadic) const override {
     const llvm::Triple &T = Context.getTargetInfo().getTriple();
     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
   // only big enough to have a vtable pointer and nothing more (or less).
   bool isNearlyEmpty(const CXXRecordDecl *RD) const override {
 
     // Check that the class has a vtable pointer.
     if (!RD->isDynamicClass())

clang/lib/AST/MicrosoftCXXABI.cpp

public:
 
   MemberPointerInfo
   getMemberPointerInfo(const MemberPointerType *MPT) const override;
 
   CallingConv getDefaultMethodCallConv(bool isVariadic) const override {
     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 {
     llvm_unreachable("unapplicable to the MS ABI");
   }
 
   const CXXConstructorDecl *
   getCopyConstructorForExceptionObject(CXXRecordDecl *RD) override {

clang/test/CodeGenOpenCLCXX/addrspace-derived-base.cl

 // RUN: %clang_cc1 %s -triple spir -cl-std=clc++ -emit-llvm -O0 -o - | FileCheck %s
 
 struct B {
   int mb;
 };
 
 class D : public B {
 public:
   int getmb() { return mb; }
 };
 
 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)*

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

 // EXPL: @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* %this
 // EXPL: @_ZNU3AS41CC2ERU3AS4KS_(%class.C addrspace(4)* %this
 // EXPL: @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* %this
 // EXPL: @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %this
 // EXPL: @_ZNU3AS4R1CaSEOU3AS4S_(%class.C addrspace(4)* %this
 // 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]]
 
 // 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]]
 
 // 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]]
 
 // 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)             \
   __kernel void test##AS() { \
     AS C c;                  \
     int i = c.get();         \
     C c1(c);                 \
     C c2;                    \
     c2 = c1;                 \
   }
 
 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]]
 
 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:%[.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]]

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: [[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;
   //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
   globI |= b;
   //CHECK: store i32 %add, i32 addrspace(1)* @globI
   globI += a;
   //CHECK: [[GVIV1:%[0-9]+]] = load volatile i32, i32 addrspace(1)* @globVI
   //CHECK: [[AND:%[a-z0-9]+]] = and i32 [[GVIV1]], 1
   //CHECK: store volatile i32 [[AND]], i32 addrspace(1)* @globVI
   globVI &= b;
   //CHECK: [[GVIV2:%[0-9]+]] = load volatile i32, i32 addrspace(1)* @globVI
   //CHECK: [[SUB:%[a-z0-9]+]] = sub nsw i32 [[GVIV2]], 0
   //CHECK: store volatile i32 [[SUB]], i32 addrspace(1)* @globVI
   globVI -= a;
 }
 
-//CHECK: define linkonce_odr void @_ZNU3AS41C6AssignE1E(%class.C addrspace(4)*{{[ %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]]
 //CHECK: store i32 {{%[a-z0-9]+}}, i32* [[E_ADDR]]
 //CHECK: [[THIS1:%[.a-z0-9]+]] = load %class.C addrspace(4)*, %class.C addrspace(4)** [[THIS_ADDR]]
 //CHECK: [[E:%[0-9]+]] = load i32, i32* [[E_ADDR]]
 //CHECK: [[ME:%[a-z0-9]+]] = getelementptr inbounds %class.C, %class.C addrspace(4)* [[THIS1]], i32 0, i32 0
 //CHECK: store i32 [[E]], i32 addrspace(4)* [[ME]]
 
-//CHECK: define linkonce_odr void @_ZNU3AS41C8OrAssignE1E(%class.C addrspace(4)*{{[ %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]]
 //CHECK: store i32 {{%[a-z0-9]+}}, i32* [[E_ADDR]]
 //CHECK: [[THIS1:%[.a-z0-9]+]] = load %class.C addrspace(4)*, %class.C addrspace(4)** [[THIS_ADDR]]
 //CHECK: [[E:%[0-9]+]] = load i32, i32* [[E_ADDR]]
 //CHECK: [[MI_GEP:%[a-z0-9]+]] = getelementptr inbounds %class.C, %class.C addrspace(4)* [[THIS1]], i32 0, i32 1
 //CHECK: [[MI:%[0-9]+]] = load i32, i32 addrspace(4)* [[MI_GEP]]
 //CHECK: %or = or i32 [[MI]], [[E]]

clang/test/CodeGenOpenCLCXX/addrspace-with-class.cl

 
 // CHECK: @const1 = addrspace(2) global %struct.MyType zeroinitializer
 __constant MyType const1 = 1;
 // CHECK: @const2 = addrspace(2) global %struct.MyType zeroinitializer
 __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();
 }
 
 // CHECK-LABEL: define spir_kernel void @fooLocal()
 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)

clang/test/CodeGenOpenCLCXX/method-overload-address-space.cl

 __global C c1;
 
 __kernel void k() {
   __local C c2;
   C c3;
   __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();
 }

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

 // RUN: %clang_cc1 -cl-std=clc++ %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:  [[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;
   S<int*> sintptr;
   S<__global int*> sintptrgl;
 
   sint.foo();
   sintptr.foo();
   sintptrgl.foo();
 }