Index: lib/AST/DeclCXX.cpp =================================================================== --- lib/AST/DeclCXX.cpp +++ lib/AST/DeclCXX.cpp @@ -2185,6 +2185,13 @@ QualType ClassTy = C.getTypeDeclType(getParent()); ClassTy = C.getQualifiedType(ClassTy, Qualifiers::fromCVRUMask(getTypeQualifiers())); + + // For OpenCL add generic address space for 'this' pointer. This allows to + // instantiate classes in various address spaces (except for constant) to be + // used with the same version of generated operators. + if (C.getLangOpts().OpenCLCPlusPlus) + ClassTy = C.getAddrSpaceQualType(ClassTy, LangAS::opencl_generic); + return C.getPointerType(ClassTy); } Index: lib/CodeGen/CGCall.cpp =================================================================== --- lib/CodeGen/CGCall.cpp +++ lib/CodeGen/CGCall.cpp @@ -72,6 +72,13 @@ /// FIXME: address space qualification? static CanQualType GetThisType(ASTContext &Context, const CXXRecordDecl *RD) { QualType RecTy = Context.getTagDeclType(RD)->getCanonicalTypeInternal(); + + if (Context.getLangOpts().OpenCLCPlusPlus) + // For OpenCL add generic address space for 'this' pointer. This allows to + // instantiate classes in various address spaces (except for constant) to be + // used with the same version of generated operators. + RecTy = Context.getAddrSpaceQualType(RecTy, LangAS::opencl_generic); + return Context.getPointerType(CanQualType::CreateUnsafe(RecTy)); } @@ -4019,11 +4026,13 @@ V->getType()->isIntegerTy()) V = Builder.CreateZExt(V, ArgInfo.getCoerceToType()); - // If the argument doesn't match, perform a bitcast to coerce it. This - // can happen due to trivial type mismatches. + // If the argument doesn't match, perform a bitcast or addrspacecast to + // coerce it. This can happen due to trivial type mismatches. if (FirstIRArg < IRFuncTy->getNumParams() && - V->getType() != IRFuncTy->getParamType(FirstIRArg)) - V = Builder.CreateBitCast(V, IRFuncTy->getParamType(FirstIRArg)); + V->getType() != IRFuncTy->getParamType(FirstIRArg)) { + auto DestTy = IRFuncTy->getParamType(FirstIRArg); + V = Builder.CreatePointerBitCastOrAddrSpaceCast(V, DestTy); + } IRCallArgs[FirstIRArg] = V; break; Index: lib/CodeGen/CGDeclCXX.cpp =================================================================== --- lib/CodeGen/CGDeclCXX.cpp +++ lib/CodeGen/CGDeclCXX.cpp @@ -26,7 +26,10 @@ static void EmitDeclInit(CodeGenFunction &CGF, const VarDecl &D, ConstantAddress DeclPtr) { - assert(D.hasGlobalStorage() && "VarDecl must have global storage!"); + assert( + (D.hasGlobalStorage() || + (D.hasLocalStorage() && CGF.getContext().getLangOpts().OpenCLCPlusPlus)) && + "VarDecl must have global or local (in the case of OpenCL) storage!"); assert(!D.getType()->isReferenceType() && "Should not call EmitDeclInit on a reference!"); Index: test/CodeGenOpenCLCXX/addrspace-of-this.cl =================================================================== --- /dev/null +++ test/CodeGenOpenCLCXX/addrspace-of-this.cl @@ -0,0 +1,102 @@ +// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s +// expected-no-diagnostics + +// Test that the 'this' pointer is in the __generic address space. + +// FIXME: Add support for __constant address space. + +class C { +public: + int v; + C() { v = 2; } + C(const C &c) { v = c.v; } + C &operator=(const C &c) { + v = c.v; + return *this; + } + int get() { return v; } +}; + +__global C c; + +__kernel void test__global() { + int i = c.get(); + C c2 = c; + C c3(c); +} + +// CHECK-LABEL: @__cxx_global_var_init() +// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) #4 + +// Test that the address space is __generic for the constructor +// CHECK-LABEL: @_ZN1CC1Ev(%class.C addrspace(4)* %this) +// CHECK: entry: +// CHECK: %this.addr = alloca %class.C addrspace(4)*, align 4 +// CHECK: store %class.C addrspace(4)* %this, %class.C addrspace(4)** %this.addr, align 4 +// CHECK: %this1 = load %class.C addrspace(4)*, %class.C addrspace(4)** %this.addr, align 4 +// CHECK: call void @_ZN1CC2Ev(%class.C addrspace(4)* %this1) #4 +// CHECK: ret void + +// CHECK-LABEL: @_Z12test__globalv() + +// Test the address space of 'this' when accessing a method. +// CHECK: %call = call i32 @_ZN1C3getEv(%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. +// CHECK: %0 = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %0, %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 assignment operator. +// CHECK: %1 = addrspacecast %class.C* %c3 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %1, %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) + +The following checkers are testing similar things like the __global test but for other address spaces. + +#define TEST(AS) \ + __kernel void test##AS() { \ + AS C c; \ + int i = c.get(); \ + C c2 = c; \ + C c3(c); \ + } + +TEST(__local) + +// CHECK-LABEL: _Z11test__localv +// CHECK: @__cxa_guard_acquire +// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*)) +// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*)) +// CHECK: %2 = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %2, %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*)) +// CHECK: %3 = addrspacecast %class.C* %c3 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %3, %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*)) + +TEST(__private) + +// CHECK-LABEL: @_Z13test__privatev +// CHECK-NOT: @__cxa_guard_acquire +// CHECK: %0 = addrspacecast %class.C* %c to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %0) +// CHECK: %1 = addrspacecast %class.C* %c to %class.C addrspace(4)* +// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* %1) +// CHECK: %2 = addrspacecast %class.C* %c to %class.C addrspace(4)* +// CHECK: %3 = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %3, %class.C addrspace(4)* dereferenceable(4) %2) +// CHECK: %4 = addrspacecast %class.C* %c to %class.C addrspace(4)* +// CHECK: %5 = addrspacecast %class.C* %c3 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %5, %class.C addrspace(4)* dereferenceable(4) %4) + +TEST() + +// CHECK-LABEL: @_Z4testv() +// CHECK: %0 = addrspacecast %class.C* %c to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %0) +// CHECK: %1 = addrspacecast %class.C* %c to %class.C addrspace(4)* +// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* %1) +// CHECK: %2 = addrspacecast %class.C* %c to %class.C addrspace(4)* +// CHECK: %3 = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %3, %class.C addrspace(4)* dereferenceable(4) %2) +// CHECK: %4 = addrspacecast %class.C* %c to %class.C addrspace(4)* +// CHECK: %5 = addrspacecast %class.C* %c3 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %5, %class.C addrspace(4)* dereferenceable(4) %4) +