Index: include/clang/AST/CanonicalType.h =================================================================== --- include/clang/AST/CanonicalType.h +++ include/clang/AST/CanonicalType.h @@ -510,7 +510,7 @@ } LLVM_CLANG_CANPROXY_SIMPLE_ACCESSOR(bool, isVariadic) - LLVM_CLANG_CANPROXY_SIMPLE_ACCESSOR(unsigned, getTypeQuals) + LLVM_CLANG_CANPROXY_SIMPLE_ACCESSOR(Qualifiers, getTypeQuals) using param_type_iterator = CanTypeIterator; Index: include/clang/AST/DeclCXX.h =================================================================== --- include/clang/AST/DeclCXX.h +++ include/clang/AST/DeclCXX.h @@ -2182,7 +2182,7 @@ /// 'this' type. QualType getThisType(ASTContext &C) const; - unsigned getTypeQualifiers() const { + Qualifiers getTypeQualifiers() const { return getType()->getAs()->getTypeQuals(); } Index: include/clang/AST/Type.h =================================================================== --- include/clang/AST/Type.h +++ include/clang/AST/Type.h @@ -256,28 +256,24 @@ } bool hasConst() const { return Mask & Const; } - void setConst(bool flag) { - Mask = (Mask & ~Const) | (flag ? Const : 0); - } + bool hasOnlyConst() const { return Mask == Const; } void removeConst() { Mask &= ~Const; } void addConst() { Mask |= Const; } bool hasVolatile() const { return Mask & Volatile; } - void setVolatile(bool flag) { - Mask = (Mask & ~Volatile) | (flag ? Volatile : 0); - } + bool hasOnlyVolatile() const { return Mask == Volatile; } void removeVolatile() { Mask &= ~Volatile; } void addVolatile() { Mask |= Volatile; } bool hasRestrict() const { return Mask & Restrict; } - void setRestrict(bool flag) { - Mask = (Mask & ~Restrict) | (flag ? Restrict : 0); - } + bool hasOnlyRestrict() const { return Mask == Restrict; } void removeRestrict() { Mask &= ~Restrict; } void addRestrict() { Mask |= Restrict; } 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; @@ -1526,7 +1522,9 @@ /// /// C++ 8.3.5p4: The return type, the parameter type list and the /// cv-qualifier-seq, [...], are part of the function type. - unsigned TypeQuals : 4; + unsigned CVRTypeQuals : 3; + /// Whether this function has extended Qualifiers. + unsigned HasExtQuals : 1; /// The number of parameters this function has, not counting '...'. /// According to [implimits] 8 bits should be enough here but this is @@ -3611,7 +3609,9 @@ FunctionTypeBits.ExtInfo = Info.Bits; } - unsigned getTypeQuals() const { return FunctionTypeBits.TypeQuals; } + Qualifiers getCVRTypeQuals() const { + return Qualifiers::fromCVRMask(FunctionTypeBits.CVRTypeQuals); + } public: QualType getReturnType() const { return ResultType; } @@ -3626,9 +3626,10 @@ CallingConv getCallConv() const { return getExtInfo().getCC(); } ExtInfo getExtInfo() const { return ExtInfo(FunctionTypeBits.ExtInfo); } - bool isConst() const { return getTypeQuals() & Qualifiers::Const; } - bool isVolatile() const { return getTypeQuals() & Qualifiers::Volatile; } - bool isRestrict() const { return getTypeQuals() & Qualifiers::Restrict; } + + bool isConst() const { return getCVRTypeQuals().hasConst(); } + bool isVolatile() const { return getCVRTypeQuals().hasVolatile(); } + bool isRestrict() const { return getCVRTypeQuals().hasRestrict(); } /// Determine the type of an expression that calls a function of /// this type. @@ -3689,7 +3690,7 @@ private llvm::TrailingObjects< FunctionProtoType, QualType, FunctionType::FunctionTypeExtraBitfields, FunctionType::ExceptionType, Expr *, FunctionDecl *, - FunctionType::ExtParameterInfo> { + FunctionType::ExtParameterInfo, Qualifiers> { friend class ASTContext; // ASTContext creates these. friend TrailingObjects; @@ -3717,6 +3718,10 @@ // an ExtParameterInfo for each of the parameters. Present if and // only if hasExtParameterInfos() is true. // + // * Optionally a Qualifiers object to represent extra qualifiers that can't + // be represented by FunctionTypeBitfields.CVRTypeQuals. Present if and only + // if hasExtQualifiers() is true. + // // The optional FunctionTypeExtraBitfields has to be before the data // related to the exception specification since it contains the number // of exception types. @@ -3763,7 +3768,7 @@ FunctionType::ExtInfo ExtInfo; bool Variadic : 1; bool HasTrailingReturn : 1; - unsigned char TypeQuals = 0; + Qualifiers TypeQuals; RefQualifierKind RefQualifier = RQ_None; ExceptionSpecInfo ExceptionSpec; const ExtParameterInfo *ExtParameterInfos = nullptr; @@ -3875,6 +3880,10 @@ return hasExtraBitfields(getExceptionSpecType()); } + bool hasExtQualifiers() const { + return FunctionTypeBits.HasExtQuals; + } + public: unsigned getNumParams() const { return FunctionTypeBits.NumParams; } @@ -3893,7 +3902,7 @@ EPI.Variadic = isVariadic(); EPI.HasTrailingReturn = hasTrailingReturn(); EPI.ExceptionSpec.Type = getExceptionSpecType(); - EPI.TypeQuals = static_cast(getTypeQuals()); + EPI.TypeQuals = getTypeQuals(); EPI.RefQualifier = getRefQualifier(); if (EPI.ExceptionSpec.Type == EST_Dynamic) { EPI.ExceptionSpec.Exceptions = exceptions(); @@ -4003,7 +4012,12 @@ /// Whether this function prototype has a trailing return type. bool hasTrailingReturn() const { return FunctionTypeBits.HasTrailingReturn; } - unsigned getTypeQuals() const { return FunctionType::getTypeQuals(); } + Qualifiers getTypeQuals() const { + if (hasExtQualifiers()) + return *getTrailingObjects(); + else + return getCVRTypeQuals(); + } /// Retrieve the ref-qualifier associated with this function type. RefQualifierKind getRefQualifier() const { Index: lib/AST/ASTContext.cpp =================================================================== --- lib/AST/ASTContext.cpp +++ lib/AST/ASTContext.cpp @@ -3768,10 +3768,11 @@ size_t Size = FunctionProtoType::totalSizeToAlloc< QualType, FunctionType::FunctionTypeExtraBitfields, FunctionType::ExceptionType, Expr *, FunctionDecl *, - FunctionProtoType::ExtParameterInfo>( + FunctionProtoType::ExtParameterInfo, Qualifiers>( NumArgs, FunctionProtoType::hasExtraBitfields(EPI.ExceptionSpec.Type), ESH.NumExceptionType, ESH.NumExprPtr, ESH.NumFunctionDeclPtr, - EPI.ExtParameterInfos ? NumArgs : 0); + EPI.ExtParameterInfos ? NumArgs : 0, + EPI.TypeQuals.hasNonFastQualifiers() ? 1 : 0); auto *FTP = (FunctionProtoType *)Allocate(Size, TypeAlignment); FunctionProtoType::ExtProtoInfo newEPI = EPI; Index: lib/AST/ASTDumper.cpp =================================================================== --- lib/AST/ASTDumper.cpp +++ lib/AST/ASTDumper.cpp @@ -337,6 +337,13 @@ if (T->isConst()) OS << " const"; if (T->isVolatile()) OS << " volatile"; if (T->isRestrict()) OS << " restrict"; + switch (T->getTypeQuals().getAddressSpace()) { + case LangAS::opencl_generic: + OS << " __generic"; + break; + default: + break; + } switch (EPI.RefQualifier) { case RQ_None: break; case RQ_LValue: OS << " &"; break; 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(); @@ -2724,7 +2724,7 @@ // Mangle CV-qualifiers, if present. These are 'this' qualifiers, // e.g. "const" in "int (A::*)() const". - mangleQualifiers(Qualifiers::fromCVRUMask(T->getTypeQuals())); + mangleQualifiers(T->getTypeQuals()); // Mangle instantiation-dependent exception-specification, if present, // per cxx-abi-dev proposal on 2016-10-11. Index: lib/AST/MicrosoftMangle.cpp =================================================================== --- lib/AST/MicrosoftMangle.cpp +++ lib/AST/MicrosoftMangle.cpp @@ -2093,7 +2093,7 @@ // If this is a C++ instance method, mangle the CVR qualifiers for the // this pointer. if (HasThisQuals) { - Qualifiers Quals = Qualifiers::fromCVRUMask(Proto->getTypeQuals()); + Qualifiers Quals = Proto->getTypeQuals(); manglePointerExtQualifiers(Quals, /*PointeeType=*/QualType()); mangleRefQualifier(Proto->getRefQualifier()); mangleQualifiers(Quals, /*IsMember=*/false); Index: lib/AST/Type.cpp =================================================================== --- lib/AST/Type.cpp +++ lib/AST/Type.cpp @@ -2842,7 +2842,7 @@ result->isInstantiationDependentType(), result->isVariablyModifiedType(), result->containsUnexpandedParameterPack(), epi.ExtInfo) { - FunctionTypeBits.TypeQuals = epi.TypeQuals; + FunctionTypeBits.CVRTypeQuals = epi.TypeQuals.getCVRQualifiers(); FunctionTypeBits.RefQualifier = epi.RefQualifier; FunctionTypeBits.NumParams = params.size(); assert(getNumParams() == params.size() && "NumParams overflow!"); @@ -2941,6 +2941,13 @@ for (unsigned i = 0; i != getNumParams(); ++i) extParamInfos[i] = epi.ExtParameterInfos[i]; } + + if (epi.TypeQuals.hasNonFastQualifiers()) { + FunctionTypeBits.HasExtQuals = 1; + *getTrailingObjects() = epi.TypeQuals; + } else { + FunctionTypeBits.HasExtQuals = 0; + } } bool FunctionProtoType::hasDependentExceptionSpec() const { @@ -3032,14 +3039,13 @@ // shortcut, use one AddInteger call instead of four for the next four // fields. assert(!(unsigned(epi.Variadic) & ~1) && - !(unsigned(epi.TypeQuals) & ~255) && !(unsigned(epi.RefQualifier) & ~3) && !(unsigned(epi.ExceptionSpec.Type) & ~15) && "Values larger than expected."); ID.AddInteger(unsigned(epi.Variadic) + - (epi.TypeQuals << 1) + - (epi.RefQualifier << 9) + - (epi.ExceptionSpec.Type << 11)); + (epi.RefQualifier << 1) + + (epi.ExceptionSpec.Type << 3)); + ID.Add(epi.TypeQuals); if (epi.ExceptionSpec.Type == EST_Dynamic) { for (QualType Ex : epi.ExceptionSpec.Exceptions) ID.AddPointer(Ex.getAsOpaquePtr()); Index: lib/AST/TypePrinter.cpp =================================================================== --- lib/AST/TypePrinter.cpp +++ lib/AST/TypePrinter.cpp @@ -801,7 +801,7 @@ printFunctionAfter(Info, OS); - if (unsigned quals = T->getTypeQuals()) { + if (unsigned quals = T->getTypeQuals().getCVRUQualifiers()) { OS << ' '; AppendTypeQualList(OS, quals, Policy.Restrict); } 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/CGDebugInfo.cpp =================================================================== --- lib/CodeGen/CGDebugInfo.cpp +++ lib/CodeGen/CGDebugInfo.cpp @@ -2590,7 +2590,7 @@ Ty->getPointeeType()->getAs(); return DBuilder.createMemberPointerType( getOrCreateInstanceMethodType(CGM.getContext().getPointerType(QualType( - Ty->getClass(), FPT->getTypeQuals())), + Ty->getClass(), FPT->getTypeQuals().getCVRUQualifiers())), FPT, U), ClassType, Size, /*Align=*/0, Flags); } 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/CGExpr.cpp =================================================================== --- lib/CodeGen/CGExpr.cpp +++ lib/CodeGen/CGExpr.cpp @@ -3943,7 +3943,7 @@ LValue RefLVal = MakeAddrLValue(addr, FieldType, FieldBaseInfo, FieldTBAAInfo); if (RecordCVR & Qualifiers::Volatile) - RefLVal.getQuals().setVolatile(true); + RefLVal.getQuals().addVolatile(); addr = EmitLoadOfReference(RefLVal, &FieldBaseInfo, &FieldTBAAInfo); // Qualifiers on the struct don't apply to the referencee. Index: lib/CodeGen/CGValue.h =================================================================== --- lib/CodeGen/CGValue.h +++ lib/CodeGen/CGValue.h @@ -562,7 +562,10 @@ } void setVolatile(bool flag) { - Quals.setVolatile(flag); + if (flag) + Quals.addVolatile(); + else + Quals.removeVolatile(); } Qualifiers::ObjCLifetime getObjCLifetime() const { 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) { @@ -2733,17 +2733,17 @@ // FIXME: Add ref-qualifier! // Handle single qualifiers without copying - if (Proto->getTypeQuals() == Qualifiers::Const) { + if (Proto->getTypeQuals().hasOnlyConst()) { Result.AddInformativeChunk(" const"); return; } - if (Proto->getTypeQuals() == Qualifiers::Volatile) { + if (Proto->getTypeQuals().hasOnlyVolatile()) { Result.AddInformativeChunk(" volatile"); return; } - if (Proto->getTypeQuals() == Qualifiers::Restrict) { + if (Proto->getTypeQuals().hasOnlyRestrict()) { Result.AddInformativeChunk(" restrict"); return; } @@ -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()); @@ -9994,7 +9994,7 @@ const FunctionProtoType *FPT = MD->getType()->castAs(); FunctionProtoType::ExtProtoInfo EPI = FPT->getExtProtoInfo(); - EPI.TypeQuals |= Qualifiers::Const; + EPI.TypeQuals.addConst(); MD->setType(Context.getFunctionType(FPT->getReturnType(), FPT->getParamTypes(), EPI)); Index: lib/Sema/SemaDeclCXX.cpp =================================================================== --- lib/Sema/SemaDeclCXX.cpp +++ lib/Sema/SemaDeclCXX.cpp @@ -8165,7 +8165,7 @@ return R; FunctionProtoType::ExtProtoInfo EPI = Proto->getExtProtoInfo(); - EPI.TypeQuals = 0; + EPI.TypeQuals = Qualifiers(); EPI.RefQualifier = RQ_None; return Context.getFunctionType(Context.VoidTy, Proto->getParamTypes(), EPI); @@ -8371,7 +8371,7 @@ const FunctionProtoType *Proto = R->getAs(); FunctionProtoType::ExtProtoInfo EPI = Proto->getExtProtoInfo(); EPI.Variadic = false; - EPI.TypeQuals = 0; + EPI.TypeQuals = Qualifiers(); EPI.RefQualifier = RQ_None; return Context.getFunctionType(Context.VoidTy, None, EPI); } @@ -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/SemaExpr.cpp =================================================================== --- lib/Sema/SemaExpr.cpp +++ lib/Sema/SemaExpr.cpp @@ -13426,7 +13426,7 @@ // Drop the parameters. FunctionProtoType::ExtProtoInfo EPI; EPI.HasTrailingReturn = false; - EPI.TypeQuals |= DeclSpec::TQ_const; + EPI.TypeQuals.addConst(); T = Context.getFunctionType(Context.DependentTy, None, EPI); Sig = Context.getTrivialTypeSourceInfo(T); } @@ -13602,7 +13602,7 @@ } else { const FunctionProtoType *FPT = cast(FTy); FunctionProtoType::ExtProtoInfo EPI = FPT->getExtProtoInfo(); - EPI.TypeQuals = 0; // FIXME: silently? + EPI.TypeQuals = Qualifiers(); EPI.ExtInfo = Ext; BlockTy = Context.getFunctionType(RetTy, FPT->getParamTypes(), EPI); } Index: lib/Sema/SemaLambda.cpp =================================================================== --- lib/Sema/SemaLambda.cpp +++ lib/Sema/SemaLambda.cpp @@ -859,7 +859,7 @@ FunctionProtoType::ExtProtoInfo EPI(Context.getDefaultCallingConvention( /*IsVariadic=*/false, /*IsCXXMethod=*/true)); EPI.HasTrailingReturn = true; - EPI.TypeQuals |= DeclSpec::TQ_const; + EPI.TypeQuals.addConst(); // C++1y [expr.prim.lambda]: // The lambda return type is 'auto', which is replaced by the // trailing-return type if provided and/or deduced from 'return' @@ -1198,7 +1198,7 @@ CallingConv CC = Context.getDefaultCallingConvention( CallOpProto->isVariadic(), /*IsCXXMethod=*/false); InvokerExtInfo.ExtInfo = InvokerExtInfo.ExtInfo.withCallingConv(CC); - InvokerExtInfo.TypeQuals = 0; + InvokerExtInfo.TypeQuals = Qualifiers(); assert(InvokerExtInfo.RefQualifier == RQ_None && "Lambda's call operator should not have a reference qualifier"); return Context.getFunctionType(CallOpProto->getReturnType(), @@ -1229,7 +1229,8 @@ S.Context.getDefaultCallingConvention( /*IsVariadic=*/false, /*IsCXXMethod=*/true)); // The conversion function is always const. - ConvExtInfo.TypeQuals = Qualifiers::Const; + ConvExtInfo.TypeQuals = Qualifiers(); + ConvExtInfo.TypeQuals.addConst(); QualType ConvTy = S.Context.getFunctionType(PtrToFunctionTy, None, ConvExtInfo); @@ -1377,7 +1378,8 @@ FunctionProtoType::ExtProtoInfo ConversionEPI( S.Context.getDefaultCallingConvention( /*IsVariadic=*/false, /*IsCXXMethod=*/true)); - ConversionEPI.TypeQuals = Qualifiers::Const; + ConversionEPI.TypeQuals = Qualifiers(); + ConversionEPI.TypeQuals.addConst(); QualType ConvTy = S.Context.getFunctionType(BlockPtrTy, None, ConversionEPI); SourceLocation Loc = IntroducerRange.getBegin(); 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; @@ -2823,8 +2823,8 @@ return; } - unsigned FromQuals = FromFunction->getTypeQuals(), - ToQuals = ToFunction->getTypeQuals(); + unsigned FromQuals = FromFunction->getTypeQuals().getCVRUQualifiers(); + unsigned ToQuals = ToFunction->getTypeQuals().getCVRUQualifiers(); if (FromQuals != ToQuals) { PDiag << ft_qualifer_mismatch << ToQuals << FromQuals; return; @@ -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; } @@ -12816,7 +12822,7 @@ // Check that the object type isn't more qualified than the // member function we're calling. - Qualifiers funcQuals = Qualifiers::fromCVRMask(proto->getTypeQuals()); + Qualifiers funcQuals = proto->getTypeQuals(); QualType objectType = op->getLHS()->getType(); if (op->getOpcode() == BO_PtrMemI) Index: lib/Sema/SemaTemplate.cpp =================================================================== --- lib/Sema/SemaTemplate.cpp +++ lib/Sema/SemaTemplate.cpp @@ -8107,7 +8107,7 @@ if (OldMD && OldMD->isConst()) { const FunctionProtoType *FPT = FT->castAs(); FunctionProtoType::ExtProtoInfo EPI = FPT->getExtProtoInfo(); - EPI.TypeQuals |= Qualifiers::Const; + EPI.TypeQuals.addConst(); FT = Context.getFunctionType(FPT->getReturnType(), FPT->getParamTypes(), EPI); } 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 @@ -1864,8 +1864,7 @@ } static std::string getFunctionQualifiersAsString(const FunctionProtoType *FnTy){ - std::string Quals = - Qualifiers::fromCVRMask(FnTy->getTypeQuals()).getAsString(); + std::string Quals = FnTy->getTypeQuals().getAsString(); switch (FnTy->getRefQualifier()) { case RQ_None: @@ -1907,7 +1906,7 @@ QualifiedFunctionKind QFK) { // Does T refer to a function type with a cv-qualifier or a ref-qualifier? const FunctionProtoType *FPT = T->getAs(); - if (!FPT || (FPT->getTypeQuals() == 0 && FPT->getRefQualifier() == RQ_None)) + if (!FPT || (FPT->getTypeQuals().empty() && FPT->getRefQualifier() == RQ_None)) return false; S.Diag(Loc, diag::err_compound_qualified_function_type) @@ -3940,7 +3939,7 @@ // Does T refer to a function type with a cv-qualifier or a ref-qualifier? bool IsQualifiedFunction = T->isFunctionProtoType() && - (T->castAs()->getTypeQuals() != 0 || + (!T->castAs()->getTypeQuals().empty() || T->castAs()->getRefQualifier() != RQ_None); // If T is 'decltype(auto)', the only declarators we can have are parens @@ -4686,7 +4685,7 @@ EPI.ExtInfo = EI; EPI.Variadic = FTI.isVariadic; EPI.HasTrailingReturn = FTI.hasTrailingReturnType(); - EPI.TypeQuals = FTI.TypeQuals; + EPI.TypeQuals.addCVRUQualifiers(FTI.TypeQuals); EPI.RefQualifier = !FTI.hasRefQualifier()? RQ_None : FTI.RefQualifierIsLValueRef? RQ_LValue : RQ_RValue; @@ -4813,7 +4812,21 @@ Exceptions, EPI.ExceptionSpec); - T = Context.getFunctionType(T, ParamTys, EPI); + // OpenCLCPlusPlus: A class member function has an address space. + if (state.getSema().getLangOpts().OpenCLCPlusPlus && + state.getDeclarator().getContext() == + DeclaratorContext::MemberContext) { + LangAS CurAS = EPI.TypeQuals.getAddressSpace(); + // If a class member function's address space is not set, set it to + // __generic. + LangAS AS = + (CurAS == LangAS::Default ? LangAS::opencl_generic : CurAS); + EPI.TypeQuals.addAddressSpace(AS); + T = Context.getFunctionType(T, ParamTys, EPI); + T = state.getSema().Context.getAddrSpaceQualType(T, AS); + } else { + T = Context.getFunctionType(T, ParamTys, EPI); + } } break; } @@ -5004,7 +5017,7 @@ // Strip the cv-qualifiers and ref-qualifiers from the type. FunctionProtoType::ExtProtoInfo EPI = FnTy->getExtProtoInfo(); - EPI.TypeQuals = 0; + EPI.TypeQuals.removeCVRQualifiers(); EPI.RefQualifier = RQ_None; T = Context.getFunctionType(FnTy->getReturnType(), FnTy->getParamTypes(), Index: lib/Serialization/ASTReader.cpp =================================================================== --- lib/Serialization/ASTReader.cpp +++ lib/Serialization/ASTReader.cpp @@ -6051,7 +6051,7 @@ EPI.Variadic = Record[Idx++]; EPI.HasTrailingReturn = Record[Idx++]; - EPI.TypeQuals = Record[Idx++]; + EPI.TypeQuals = Qualifiers::fromOpaqueValue(Record[Idx++]); EPI.RefQualifier = static_cast(Record[Idx++]); SmallVector ExceptionStorage; readExceptionSpec(*Loc.F, ExceptionStorage, EPI.ExceptionSpec, Record, Idx); Index: lib/Serialization/ASTWriter.cpp =================================================================== --- lib/Serialization/ASTWriter.cpp +++ lib/Serialization/ASTWriter.cpp @@ -310,7 +310,7 @@ Record.push_back(T->isVariadic()); Record.push_back(T->hasTrailingReturn()); - Record.push_back(T->getTypeQuals()); + Record.push_back(T->getTypeQuals().getAsOpaqueValue()); Record.push_back(static_cast(T->getRefQualifier())); addExceptionSpec(T, Record); 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) {