Index: lib/AST/DeclCXX.cpp =================================================================== --- lib/AST/DeclCXX.cpp +++ lib/AST/DeclCXX.cpp @@ -2185,6 +2185,8 @@ QualType ClassTy = C.getTypeDeclType(getParent()); ClassTy = C.getQualifiedType(ClassTy, Qualifiers::fromCVRUMask(getTypeQualifiers())); + + ClassTy = C.getAddrSpaceQualType(ClassTy, getType().getAddressSpace()); return C.getPointerType(ClassTy); } Index: lib/CodeGen/CGCall.cpp =================================================================== --- lib/CodeGen/CGCall.cpp +++ lib/CodeGen/CGCall.cpp @@ -69,9 +69,10 @@ /// Derives the 'this' type for codegen purposes, i.e. ignoring method /// qualification. -/// FIXME: address space qualification? -static CanQualType GetThisType(ASTContext &Context, const CXXRecordDecl *RD) { +static CanQualType GetThisType(ASTContext &Context, const CXXRecordDecl *RD, const CXXMethodDecl *MD) { QualType RecTy = Context.getTagDeclType(RD)->getCanonicalTypeInternal(); + if (MD) + RecTy = Context.getAddrSpaceQualType(RecTy, MD->getType().getAddressSpace()); return Context.getPointerType(CanQualType::CreateUnsafe(RecTy)); } @@ -246,7 +247,7 @@ // Add the 'this' pointer. if (RD) - argTypes.push_back(GetThisType(Context, RD)); + argTypes.push_back(GetThisType(Context, RD, MD)); else argTypes.push_back(Context.VoidPtrTy); @@ -302,7 +303,7 @@ SmallVector argTypes; SmallVector paramInfos; - argTypes.push_back(GetThisType(Context, MD->getParent())); + argTypes.push_back(GetThisType(Context, MD->getParent(), MD)); bool PassParams = true; @@ -529,7 +530,7 @@ CodeGenTypes::arrangeUnprototypedMustTailThunk(const CXXMethodDecl *MD) { assert(MD->isVirtual() && "only methods have thunks"); CanQual FTP = GetFormalType(MD); - CanQualType ArgTys[] = { GetThisType(Context, MD->getParent()) }; + CanQualType ArgTys[] = { GetThisType(Context, MD->getParent(), MD) }; return arrangeLLVMFunctionInfo(Context.VoidTy, /*instanceMethod=*/false, /*chainCall=*/false, ArgTys, FTP->getExtInfo(), {}, RequiredArgs(1)); @@ -543,7 +544,7 @@ CanQual FTP = GetFormalType(CD); SmallVector ArgTys; const CXXRecordDecl *RD = CD->getParent(); - ArgTys.push_back(GetThisType(Context, RD)); + ArgTys.push_back(GetThisType(Context, RD, CD)); if (CT == Ctor_CopyingClosure) ArgTys.push_back(*FTP->param_type_begin()); if (RD->getNumVBases() > 0) Index: lib/CodeGen/CGClass.cpp =================================================================== --- lib/CodeGen/CGClass.cpp +++ lib/CodeGen/CGClass.cpp @@ -16,6 +16,7 @@ #include "CGDebugInfo.h" #include "CGRecordLayout.h" #include "CodeGenFunction.h" +#include "TargetInfo.h" #include "clang/AST/CXXInheritance.h" #include "clang/AST/DeclTemplate.h" #include "clang/AST/EvaluatedExprVisitor.h" @@ -2012,8 +2013,18 @@ bool NewPointerIsChecked) { CallArgList Args; + LangAS AS = D->getType().getAddressSpace(); + llvm::Value *ThisPtr = This.getPointer(); + if (AS != LangAS::Default) { + unsigned TargetAS = getContext().getTargetAddressSpace(AS); + llvm::Type *NewType = ThisPtr->getType()->getPointerTo(TargetAS); + ThisPtr = getTargetHooks().performAddrSpaceCast( + *this, This.getPointer(), + getLangASFromTargetAS(This.getAddressSpace()), AS, + NewType); + } // Push the this ptr. - Args.add(RValue::get(This.getPointer()), D->getThisType(getContext())); + Args.add(RValue::get(ThisPtr), D->getThisType(getContext())); // If this is a trivial constructor, emit a memcpy now before we lose // the alignment information on the argument. 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: lib/CodeGen/CGExprCXX.cpp =================================================================== --- lib/CodeGen/CGExprCXX.cpp +++ lib/CodeGen/CGExprCXX.cpp @@ -11,12 +11,13 @@ // //===----------------------------------------------------------------------===// -#include "CodeGenFunction.h" #include "CGCUDARuntime.h" #include "CGCXXABI.h" #include "CGDebugInfo.h" #include "CGObjCRuntime.h" +#include "CodeGenFunction.h" #include "ConstantEmitter.h" +#include "TargetInfo.h" #include "clang/CodeGen/CGFunctionInfo.h" #include "clang/Frontend/CodeGenOptions.h" #include "llvm/IR/CallSite.h" @@ -44,11 +45,26 @@ "Trying to emit a member or operator call expr on a static method!"); ASTContext &C = CGF.getContext(); - // Push the this ptr. const CXXRecordDecl *RD = CGF.CGM.getCXXABI().getThisArgumentTypeForMethod(MD); - Args.add(RValue::get(This), - RD ? C.getPointerType(C.getTypeDeclType(RD)) : C.VoidPtrTy); + + if (RD) { + LangAS AS = MD->getType().getAddressSpace(); + if (AS != LangAS::Default) { + unsigned TargetAS = C.getTargetAddressSpace(AS); + llvm::Type *NewType = This->getType()->getPointerTo(TargetAS); + This = CGF.getTargetHooks().performAddrSpaceCast( + CGF, This, + getLangASFromTargetAS(This->getType()->getPointerAddressSpace()), AS, + NewType); + } + // Push the this ptr. + Args.add(RValue::get(This), C.getPointerType(C.getAddrSpaceQualType( + C.getTypeDeclType(RD), AS))); + } else { + // Push the this ptr. + Args.add(RValue::get(This), C.VoidPtrTy); + } // If there is an implicit parameter (e.g. VTT), emit it. if (ImplicitParam) { Index: lib/Sema/SemaType.cpp =================================================================== --- lib/Sema/SemaType.cpp +++ lib/Sema/SemaType.cpp @@ -7189,12 +7189,16 @@ bool IsFuncType = ChunkIndex < D.getNumTypeObjects() && D.getTypeObject(ChunkIndex).Kind == DeclaratorChunk::Function; + bool IsClassMemberFn = IsFuncType && D.getContext() == DeclaratorContext::MemberContext; + if ( // Do not deduce addr space for function return type and function type, - // otherwise it will fail some sema check. - IsFuncReturnType || IsFuncType || + // otherwise it will fail some sema check. We want to deduce class member functions. + IsFuncReturnType || + (IsFuncType && !IsClassMemberFn) || // Do not deduce addr space for member types of struct, except the pointee - // type of a pointer member type. - (D.getContext() == DeclaratorContext::MemberContext && !IsPointee) || + // type of a pointer member type. We want to deduce class member functions. + (D.getContext() == DeclaratorContext::MemberContext && !IsPointee && + !IsClassMemberFn) || // Do not deduce addr space for types used to define a typedef and the // typedef itself, except the pointee type of a pointer type which is used // to define the typedef. @@ -7224,7 +7228,8 @@ // (...) // Pointers that are declared without pointing to a named address space // point to the generic address space. - if (IsPointee) { + // Deduce class members functions to be of the generic address space + if (IsPointee || IsClassMemberFn) { ImpAddr = LangAS::opencl_generic; } else { if (D.getContext() == DeclaratorContext::FileContext) { Index: test/CodeGenOpenCLCXX/addrspace-of-this.cl =================================================================== --- /dev/null +++ test/CodeGenOpenCLCXX/addrspace-of-this.cl @@ -0,0 +1,140 @@ +// 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 c1(c); + C c2; + c2 = c1; +} + +// 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 invoking 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* %c1 to %class.C* addrspace(4)* +// CHECK: %1 = bitcast %class.C* addrspace(4)* %0 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)*)) + +// Test the address space of 'this' when invoking a constructor. +// CHECK: %2 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)* +// CHECK: %3 = bitcast %class.C* addrspace(4)* %2 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %3) #4 + +// Test the address space of 'this' when invoking assignment operator. +// CHECK: %4 = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// CHECK: %5 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)* +// CHECK: %6 = bitcast %class.C* addrspace(4)* %5 to %class.C addrspace(4)* +// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZN1CaSERU3AS4KS_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %4) + +#define TEST(AS) \ + __kernel void test##AS() { \ + AS C c; \ + int i = c.get(); \ + C c1(c); \ + C c2; \ + c2 = c1; \ + } + +TEST(__local) + +// CHECK-LABEL: _Z11test__localv +// CHECK: @__cxa_guard_acquire + +// Test the address space of 'this' when invoking a method. +// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*)) + +// Test the address space of 'this' when invoking copy-constructor. +// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*)) + +// Test the address space of 'this' when invoking a constructor. +// CHECK: %4 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)* +// CHECK: %5 = bitcast %class.C* addrspace(4)* %4 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %5) + +// Test the address space of 'this' when invoking assignment operator. +// CHECK: %6 = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// CHECK: %7 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)* +// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZN1CaSERU3AS4KS_(%class.C addrspace(4)* %8, %class.C addrspace(4)* dereferenceable(4) %6) + +TEST(__private) + +// CHECK-LABEL: @_Z13test__privatev + +// Test the address space of 'this' when invoking a method. +// CHECK: %2 = addrspacecast %class.C* %c to %class.C* addrspace(4)* +// CHECK: %3 = bitcast %class.C* addrspace(4)* %2 to %class.C addrspace(4)* +// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* %3) + +// Test the address space of 'this' when invoking a copy-constructor. +// CHECK: %5 = addrspacecast %class.C* %c to %class.C addrspace(4)* +// CHECK: %6 = bitcast %class.C* addrspace(4)* %4 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %5) + +// Test the address space of 'this' when invoking a constructor. +// CHECK: %7 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)* +// CHECK: %8 = bitcast %class.C* addrspace(4)* %7 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %8) + +// Test the address space of 'this' when invoking a copy-assignment. +// CHECK: %9 = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// CHECK: %10 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)* +// CHECK: %11 = bitcast %class.C* addrspace(4)* %10 to %class.C addrspace(4)* +// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZN1CaSERU3AS4KS_(%class.C addrspace(4)* %11, %class.C addrspace(4)* dereferenceable(4) %9) + +TEST() + +// CHECK-LABEL: @_Z4testv() +// Test the address space of 'this' when invoking a method. +// CHECK: %2 = addrspacecast %class.C* %c to %class.C* addrspace(4)* +// CHECK: %3 = bitcast %class.C* addrspace(4)* %2 to %class.C addrspace(4)* +// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* %3) #4 + +// Test the address space of 'this' when invoking a copy-constructor. +// CHECK: %4 = addrspacecast %class.C* %c1 to %class.C* addrspace(4)* +// CHECK: %5 = addrspacecast %class.C* %c to %class.C addrspace(4)* +// CHECK: %6 = bitcast %class.C* addrspace(4)* %4 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %5) + +// Test the address space of 'this' when invoking a constructor. +// CHECK: %7 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)* +// CHECK: %8 = bitcast %class.C* addrspace(4)* %7 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %8) + +// Test the address space of 'this' when invoking a copy-assignment. +// CHECK: %9 = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// CHECK: %10 = addrspacecast %class.C* %c2 to %class.C* addrspace(4)* +// CHECK: %11 = bitcast %class.C* addrspace(4)* %10 to %class.C addrspace(4)* +// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZN1CaSERU3AS4KS_(%class.C addrspace(4)* %11, %class.C addrspace(4)* dereferenceable(4) %9)