Index: include/clang/AST/DeclCXX.h =================================================================== --- include/clang/AST/DeclCXX.h +++ include/clang/AST/DeclCXX.h @@ -2182,8 +2182,12 @@ /// 'this' type. QualType getThisType(ASTContext &C) const; - unsigned getTypeQualifiers() const { - return getType()->getAs()->getTypeQuals(); + Qualifiers getTypeQualifiers() const { + Qualifiers C; + unsigned CVRU = getType()->getAs()->getTypeQuals(); + C.addCVRUQualifiers(CVRU); + C.addQualifiers(getType().getQualifiers()); + return C; } /// Retrieve the ref-qualifier associated with this method. Index: include/clang/AST/Type.h =================================================================== --- include/clang/AST/Type.h +++ include/clang/AST/Type.h @@ -278,6 +278,8 @@ bool hasCVRQualifiers() const { return getCVRQualifiers(); } unsigned getCVRQualifiers() const { return Mask & CVRMask; } + unsigned getCVRUQualifiers() const { return Mask & (CVRMask | UMask); } + void setCVRQualifiers(unsigned mask) { assert(!(mask & ~CVRMask) && "bitmask contains non-CVR bits"); Mask = (Mask & ~CVRMask) | mask; Index: lib/AST/DeclCXX.cpp =================================================================== --- lib/AST/DeclCXX.cpp +++ lib/AST/DeclCXX.cpp @@ -2183,8 +2183,8 @@ assert(isInstance() && "No 'this' for static methods!"); QualType ClassTy = C.getTypeDeclType(getParent()); - ClassTy = C.getQualifiedType(ClassTy, - Qualifiers::fromCVRUMask(getTypeQualifiers())); + ClassTy = C.getQualifiedType(ClassTy, getTypeQualifiers()); + return C.getPointerType(ClassTy); } Index: lib/AST/ItaniumMangle.cpp =================================================================== --- lib/AST/ItaniumMangle.cpp +++ lib/AST/ItaniumMangle.cpp @@ -1503,8 +1503,8 @@ Out << 'N'; if (const CXXMethodDecl *Method = dyn_cast(ND)) { - Qualifiers MethodQuals = - Qualifiers::fromCVRUMask(Method->getTypeQualifiers()); + Qualifiers MethodQuals = Qualifiers::fromCVRUMask( + Method->getTypeQualifiers().getCVRUQualifiers()); // We do not consider restrict a distinguishing attribute for overloading // purposes so we must not mangle it. MethodQuals.removeRestrict(); Index: lib/CodeGen/CGCall.cpp =================================================================== --- lib/CodeGen/CGCall.cpp +++ lib/CodeGen/CGCall.cpp @@ -67,11 +67,12 @@ } } -/// Derives the 'this' type for codegen purposes, i.e. ignoring method +/// Derives the 'this' type for codegen purposes, i.e. ignoring method CVR /// 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,19 @@ bool NewPointerIsChecked) { CallArgList Args; + LangAS SlotAS = E->getType().getAddressSpace(); + QualType ThisType = D->getThisType(getContext()); + LangAS ThisAS = ThisType.getTypePtr()->getPointeeType().getAddressSpace(); + llvm::Value *ThisPtr = This.getPointer(); + if (SlotAS != ThisAS) { + unsigned TargetThisAS = getContext().getTargetAddressSpace(ThisAS); + llvm::Type *NewType = + ThisPtr->getType()->getPointerElementType()->getPointerTo(TargetThisAS); + ThisPtr = getTargetHooks().performAddrSpaceCast(*this, This.getPointer(), + ThisAS, SlotAS, 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/Index/USRGeneration.cpp =================================================================== --- lib/Index/USRGeneration.cpp +++ lib/Index/USRGeneration.cpp @@ -270,7 +270,7 @@ if (const CXXMethodDecl *MD = dyn_cast(D)) { if (MD->isStatic()) Out << 'S'; - if (unsigned quals = MD->getTypeQualifiers()) + if (unsigned quals = MD->getTypeQualifiers().getCVRUQualifiers()) Out << (char)('0' + quals); switch (MD->getRefQualifier()) { case RQ_None: break; Index: lib/Parse/ParseCXXInlineMethods.cpp =================================================================== --- lib/Parse/ParseCXXInlineMethods.cpp +++ lib/Parse/ParseCXXInlineMethods.cpp @@ -416,7 +416,7 @@ Method = cast(LM.Method); Sema::CXXThisScopeRAII ThisScope(Actions, Method->getParent(), - Method->getTypeQualifiers(), + Method->getTypeQualifiers().getCVRUQualifiers(), getLangOpts().CPlusPlus11); // Parse the exception-specification. Index: lib/Sema/SemaCodeComplete.cpp =================================================================== --- lib/Sema/SemaCodeComplete.cpp +++ lib/Sema/SemaCodeComplete.cpp @@ -1042,8 +1042,8 @@ if (HasObjectTypeQualifiers) if (const CXXMethodDecl *Method = dyn_cast(R.Declaration)) if (Method->isInstance()) { - Qualifiers MethodQuals - = Qualifiers::fromCVRMask(Method->getTypeQualifiers()); + Qualifiers MethodQuals = Qualifiers::fromCVRMask( + Method->getTypeQualifiers().getCVRQualifiers()); if (ObjectTypeQualifiers == MethodQuals) R.Priority += CCD_ObjectQualifierMatch; else if (ObjectTypeQualifiers - MethodQuals) { @@ -3697,8 +3697,8 @@ RecordDecl *MemberCompletionRecord = nullptr; if (CXXMethodDecl *CurMethod = dyn_cast(CurContext)) { if (CurMethod->isInstance()) { - Results.setObjectTypeQualifiers( - Qualifiers::fromCVRMask(CurMethod->getTypeQualifiers())); + Results.setObjectTypeQualifiers(Qualifiers::fromCVRMask( + CurMethod->getTypeQualifiers().getCVRQualifiers())); MemberCompletionRecord = CurMethod->getParent(); } } Index: lib/Sema/SemaDecl.cpp =================================================================== --- lib/Sema/SemaDecl.cpp +++ lib/Sema/SemaDecl.cpp @@ -9986,7 +9986,7 @@ CXXMethodDecl *MD = dyn_cast(NewFD); if (!getLangOpts().CPlusPlus14 && MD && MD->isConstexpr() && !MD->isStatic() && !isa(MD) && - (MD->getTypeQualifiers() & Qualifiers::Const) == 0) { + !MD->getTypeQualifiers().hasConst()) { CXXMethodDecl *OldMD = nullptr; if (OldDecl) OldMD = dyn_cast_or_null(OldDecl->getAsFunction()); Index: lib/Sema/SemaDeclCXX.cpp =================================================================== --- lib/Sema/SemaDeclCXX.cpp +++ lib/Sema/SemaDeclCXX.cpp @@ -11956,10 +11956,12 @@ // Dereference "this". DerefBuilder DerefThis(This); - CastBuilder To(DerefThis, - Context.getCVRQualifiedType( - BaseType, CopyAssignOperator->getTypeQualifiers()), - VK_LValue, BasePath); + CastBuilder To( + DerefThis, + Context.getCVRQualifiedType( + BaseType, + CopyAssignOperator->getTypeQualifiers().getCVRUQualifiers()), + VK_LValue, BasePath); // Build the copy. StmtResult Copy = buildSingleCopyAssign(*this, Loc, BaseType, @@ -12323,10 +12325,12 @@ DerefBuilder DerefThis(This); // Implicitly cast "this" to the appropriately-qualified base type. - CastBuilder To(DerefThis, - Context.getCVRQualifiedType( - BaseType, MoveAssignOperator->getTypeQualifiers()), - VK_LValue, BasePath); + CastBuilder To( + DerefThis, + Context.getCVRQualifiedType( + BaseType, + MoveAssignOperator->getTypeQualifiers().getCVRQualifiers()), + VK_LValue, BasePath); // Build the move. StmtResult Move = buildSingleCopyAssign(*this, Loc, BaseType, Index: lib/Sema/SemaOverload.cpp =================================================================== --- lib/Sema/SemaOverload.cpp +++ lib/Sema/SemaOverload.cpp @@ -1142,8 +1142,8 @@ // function yet (because we haven't yet resolved whether this is a static // or non-static member function). Add it now, on the assumption that this // is a redeclaration of OldMethod. - unsigned OldQuals = OldMethod->getTypeQualifiers(); - unsigned NewQuals = NewMethod->getTypeQualifiers(); + unsigned OldQuals = OldMethod->getTypeQualifiers().getCVRUQualifiers(); + unsigned NewQuals = NewMethod->getTypeQualifiers().getCVRUQualifiers(); if (!getLangOpts().CPlusPlus14 && NewMethod->isConstexpr() && !isa(NewMethod)) NewQuals |= Qualifiers::Const; @@ -5065,8 +5065,9 @@ QualType ClassType = S.Context.getTypeDeclType(ActingContext); // [class.dtor]p2: A destructor can be invoked for a const, volatile or // const volatile object. - unsigned Quals = isa(Method) ? - Qualifiers::Const | Qualifiers::Volatile : Method->getTypeQualifiers(); + unsigned Quals = isa(Method) + ? Qualifiers::Const | Qualifiers::Volatile + : Method->getTypeQualifiers().getCVRUQualifiers(); QualType ImplicitParamType = S.Context.getCVRQualifiedType(ClassType, Quals); // Set up the conversion sequence as a "bad" conversion, to allow us @@ -5249,9 +5250,14 @@ From = FromRes.get(); } - if (!Context.hasSameType(From->getType(), DestType)) - From = ImpCastExprToType(From, DestType, CK_NoOp, + if (!Context.hasSameType(From->getType(), DestType)) { + if (From->getType().getAddressSpace() != DestType.getAddressSpace()) + From = ImpCastExprToType(From, DestType, CK_AddressSpaceConversion, From->getValueKind()).get(); + else + From = ImpCastExprToType(From, DestType, CK_NoOp, + From->getValueKind()).get(); + } return From; } Index: lib/Sema/SemaTemplateDeduction.cpp =================================================================== --- lib/Sema/SemaTemplateDeduction.cpp +++ lib/Sema/SemaTemplateDeduction.cpp @@ -3080,7 +3080,7 @@ CXXRecordDecl *ThisContext = nullptr; if (CXXMethodDecl *Method = dyn_cast(Function)) { ThisContext = Method->getParent(); - ThisTypeQuals = Method->getTypeQualifiers(); + ThisTypeQuals = Method->getTypeQualifiers().getCVRQualifiers();; } CXXThisScopeRAII ThisScope(*this, ThisContext, ThisTypeQuals, @@ -4655,8 +4655,9 @@ // The standard doesn't say explicitly, but we pick the appropriate kind of // reference type based on [over.match.funcs]p4. QualType ArgTy = Context.getTypeDeclType(Method->getParent()); - ArgTy = Context.getQualifiedType(ArgTy, - Qualifiers::fromCVRMask(Method->getTypeQualifiers())); + ArgTy = Context.getQualifiedType( + ArgTy, + Qualifiers::fromCVRMask(Method->getTypeQualifiers().getCVRQualifiers())); if (Method->getRefQualifier() == RQ_RValue) ArgTy = Context.getRValueReferenceType(ArgTy); else Index: lib/Sema/SemaTemplateInstantiateDecl.cpp =================================================================== --- lib/Sema/SemaTemplateInstantiateDecl.cpp +++ lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -3425,7 +3425,7 @@ unsigned ThisTypeQuals = 0; if (CXXMethodDecl *Method = dyn_cast(D)) { ThisContext = cast(Owner); - ThisTypeQuals = Method->getTypeQualifiers(); + ThisTypeQuals = Method->getTypeQualifiers().getCVRQualifiers(); } TypeSourceInfo *NewTInfo 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,129 @@ +// 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: 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 a constructor. +// CHECK: %1 = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %1) #4 + +// Test the address space of 'this' when invoking assignment operator. +// CHECK: %2 = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// CHECK: %3 = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZN1CaSERU3AS4KS_(%class.C addrspace(4)* %3, %class.C addrspace(4)* dereferenceable(4) %2) + +#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: %3 = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %3) + +// 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: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZN1CaSERU3AS4KS_(%class.C addrspace(4)* %5, %class.C addrspace(4)* dereferenceable(4) %4) + +TEST(__private) + +// CHECK-LABEL: @_Z13test__privatev + +// Test the address space of 'this' when invoking a method. +// CHECK: %1 = addrspacecast %class.C* %c to %class.C addrspace(4)* +// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* %1) + +// Test the address space of 'this' when invoking a copy-constructor. +// CHECK: %2 = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// CHECK: %3 = addrspacecast %class.C* %c to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %2, %class.C addrspace(4)* dereferenceable(4) %3) + +// Test the address space of 'this' when invoking a constructor. +// CHECK: %4 = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %4) + +// Test the address space of 'this' when invoking a copy-assignment. +// CHECK: %5 = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// CHECK: %6 = addrspacecast %class.C* %c2 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) %5) + +TEST() + +// CHECK-LABEL: @_Z4testv() +// Test the address space of 'this' when invoking a method. +// CHECK: %1 = addrspacecast %class.C* %c to %class.C addrspace(4)* +// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* %1) #4 + +// Test the address space of 'this' when invoking a copy-constructor. +// CHECK: %2 = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// CHECK: %3 = addrspacecast %class.C* %c to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %2, %class.C addrspace(4)* dereferenceable(4) %3) + +// Test the address space of 'this' when invoking a constructor. +// CHECK: %4 = addrspacecast %class.C* %c2 to %class.C addrspace(4)* +// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %4) + +// Test the address space of 'this' when invoking a copy-assignment. +// CHECK: %5 = addrspacecast %class.C* %c1 to %class.C addrspace(4)* +// CHECK: %6 = addrspacecast %class.C* %c2 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) %5) Index: tools/libclang/CIndex.cpp =================================================================== --- tools/libclang/CIndex.cpp +++ tools/libclang/CIndex.cpp @@ -8370,7 +8370,7 @@ const Decl *D = cxcursor::getCursorDecl(C); const CXXMethodDecl *Method = D ? dyn_cast_or_null(D->getAsFunction()) : nullptr; - return (Method && (Method->getTypeQualifiers() & Qualifiers::Const)) ? 1 : 0; + return (Method && Method->getTypeQualifiers().hasConst()) ? 1 : 0; } unsigned clang_CXXMethod_isDefaulted(CXCursor C) {