Index: include/clang/Sema/ParsedAttr.h =================================================================== --- include/clang/Sema/ParsedAttr.h +++ include/clang/Sema/ParsedAttr.h @@ -568,6 +568,25 @@ /// parsed attribute does not have a semantic equivalent, or would not have /// a Spelling enumeration, the value UINT_MAX is returned. unsigned getSemanticSpelling() const; + + /// If this is an OpenCL addr space attribute returns its representation + /// in LangAS, otherwise returns default addr space. + LangAS asOpenCLLangAS() const { + switch (getKind()) { + case ParsedAttr::AT_OpenCLConstantAddressSpace: + return LangAS::opencl_constant; + case ParsedAttr::AT_OpenCLGlobalAddressSpace: + return LangAS::opencl_global; + case ParsedAttr::AT_OpenCLLocalAddressSpace: + return LangAS::opencl_local; + case ParsedAttr::AT_OpenCLPrivateAddressSpace: + return LangAS::opencl_private; + case ParsedAttr::AT_OpenCLGenericAddressSpace: + return LangAS::opencl_generic; + default: + return LangAS::Default; + } + } }; class AttributePool; Index: lib/Parse/ParseDecl.cpp =================================================================== --- lib/Parse/ParseDecl.cpp +++ lib/Parse/ParseDecl.cpp @@ -6159,6 +6159,16 @@ Qualifiers Q = Qualifiers::fromCVRUMask(DS.getTypeQualifiers()); if (D.getDeclSpec().isConstexprSpecified() && !getLangOpts().CPlusPlus14) Q.addConst(); + // FIXME: Collect C++ address spaces. + if (getLangOpts().OpenCLCPlusPlus) { + for (ParsedAttr &attr : DS.getAttributes()) { + LangAS ASIdx = attr.asOpenCLLangAS(); + if (ASIdx != LangAS::Default) { + Q.addAddressSpace(ASIdx); + break; + } + } + } Sema::CXXThisScopeRAII ThisScope( Actions, dyn_cast(Actions.CurContext), Q, Index: lib/Sema/SemaOverload.cpp =================================================================== --- lib/Sema/SemaOverload.cpp +++ lib/Sema/SemaOverload.cpp @@ -1172,16 +1172,14 @@ // 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. - // FIXME: OpenCL: Need to consider address spaces - unsigned OldQuals = OldMethod->getTypeQualifiers().getCVRUQualifiers(); - unsigned NewQuals = NewMethod->getTypeQualifiers().getCVRUQualifiers(); + auto OldQuals = OldMethod->getTypeQualifiers(); + auto NewQuals = NewMethod->getTypeQualifiers(); if (!getLangOpts().CPlusPlus14 && NewMethod->isConstexpr() && !isa(NewMethod)) - NewQuals |= Qualifiers::Const; - + NewQuals.addConst(); // We do not allow overloading based off of '__restrict'. - OldQuals &= ~Qualifiers::Restrict; - NewQuals &= ~Qualifiers::Restrict; + OldQuals.removeRestrict(); + NewQuals.removeRestrict(); if (OldQuals != NewQuals) return true; } @@ -5148,6 +5146,16 @@ return ICS; } + if (FromTypeCanon.getQualifiers().hasAddressSpace()) { + Qualifiers QualsImplicitParamType = ImplicitParamType.getQualifiers(); + Qualifiers QualsFromType = FromTypeCanon.getQualifiers(); + if (!QualsImplicitParamType.isAddressSpaceSupersetOf(QualsFromType)) { + ICS.setBad(BadConversionSequence::bad_qualifiers, + FromType, ImplicitParamType); + return ICS; + } + } + // Check that we have either the same type or a derived type. It // affects the conversion rank. QualType ClassTypeCanon = S.Context.getCanonicalType(ClassType); Index: lib/Sema/SemaType.cpp =================================================================== --- lib/Sema/SemaType.cpp +++ lib/Sema/SemaType.cpp @@ -3916,6 +3916,25 @@ llvm_unreachable("unknown NullabilityKind"); } +// Diagnose whether this is a case with the multiple addr spaces. +// Returns true if this is an invalid case. +// ISO/IEC TR 18037 S5.3 (amending C99 6.7.3): "No type shall be qualified +// by qualifiers for two or more different address spaces." +static bool DiagnoseMultipleAddrSpaceAttributes(Sema &S, LangAS ASOld, + LangAS ASNew, + SourceLocation AttrLoc) { + if (ASOld != LangAS::Default) { + if (ASOld != ASNew) { + S.Diag(AttrLoc, diag::err_attribute_address_multiple_qualifiers); + return true; + } else + // Emit a warning if they are identical; it's likely unintended. + S.Diag(AttrLoc, + diag::warn_attribute_address_multiple_identical_qualifiers); + } + return false; +} + static TypeSourceInfo * GetTypeSourceInfoForDeclarator(TypeProcessingState &State, QualType T, TypeSourceInfo *ReturnTypeInfo); @@ -4823,18 +4842,35 @@ Exceptions, EPI.ExceptionSpec); - const auto &Spec = D.getCXXScopeSpec(); + // FIXME: Set address space from attrs for C++ mode here. // OpenCLCPlusPlus: A class member function has an address space. - if (state.getSema().getLangOpts().OpenCLCPlusPlus && - ((!Spec.isEmpty() && - Spec.getScopeRep()->getKind() == NestedNameSpecifier::TypeSpec) || - state.getDeclarator().getContext() == - DeclaratorContext::MemberContext)) { - LangAS CurAS = EPI.TypeQuals.getAddressSpace(); + auto IsClassMember = [&]() { + return (!state.getDeclarator().getCXXScopeSpec().isEmpty() && + state.getDeclarator() + .getCXXScopeSpec() + .getScopeRep() + ->getKind() == NestedNameSpecifier::TypeSpec) || + state.getDeclarator().getContext() == + DeclaratorContext::MemberContext; + }; + + if (state.getSema().getLangOpts().OpenCLCPlusPlus && IsClassMember()) { + LangAS ASIdx = LangAS::Default; + // Take address space attr if any and mark as invalid to avoid adding + // them later while creating QualType. + if (FTI.MethodQualifiers) + for (ParsedAttr &attr : FTI.MethodQualifiers->getAttributes()) { + LangAS ASIdxNew = attr.asOpenCLLangAS(); + if (DiagnoseMultipleAddrSpaceAttributes(S, ASIdx, ASIdxNew, + attr.getLoc())) + D.setInvalidType(true); + else + ASIdx = ASIdxNew; + } // If a class member function's address space is not set, set it to // __generic. LangAS AS = - (CurAS == LangAS::Default ? LangAS::opencl_generic : CurAS); + (ASIdx == LangAS::Default ? LangAS::opencl_generic : ASIdx); EPI.TypeQuals.addAddressSpace(AS); } T = Context.getFunctionType(T, ParamTys, EPI); @@ -5790,19 +5826,9 @@ LangAS ASIdx = getLangASFromTargetAS(static_cast(addrSpace.getZExtValue())); - // If this type is already address space qualified with a different - // address space, reject it. - // ISO/IEC TR 18037 S5.3 (amending C99 6.7.3): "No type shall be qualified - // by qualifiers for two or more different address spaces." - if (T.getAddressSpace() != LangAS::Default) { - if (T.getAddressSpace() != ASIdx) { - Diag(AttrLoc, diag::err_attribute_address_multiple_qualifiers); - return QualType(); - } else - // Emit a warning if they are identical; it's likely unintended. - Diag(AttrLoc, - diag::warn_attribute_address_multiple_identical_qualifiers); - } + if (DiagnoseMultipleAddrSpaceAttributes(*this, T.getAddressSpace(), ASIdx, + AttrLoc)) + return QualType(); return Context.getAddrSpaceQualType(T, ASIdx); } @@ -5880,34 +5906,14 @@ } } else { // The keyword-based type attributes imply which address space to use. - switch (Attr.getKind()) { - case ParsedAttr::AT_OpenCLGlobalAddressSpace: - ASIdx = LangAS::opencl_global; break; - case ParsedAttr::AT_OpenCLLocalAddressSpace: - ASIdx = LangAS::opencl_local; break; - case ParsedAttr::AT_OpenCLConstantAddressSpace: - ASIdx = LangAS::opencl_constant; break; - case ParsedAttr::AT_OpenCLGenericAddressSpace: - ASIdx = LangAS::opencl_generic; break; - case ParsedAttr::AT_OpenCLPrivateAddressSpace: - ASIdx = LangAS::opencl_private; break; - default: + ASIdx = Attr.asOpenCLLangAS(); + if (ASIdx == LangAS::Default) llvm_unreachable("Invalid address space"); - } - // If this type is already address space qualified with a different - // address space, reject it. - // ISO/IEC TR 18037 S5.3 (amending C99 6.7.3): "No type shall be qualified by - // qualifiers for two or more different address spaces." - if (Type.getAddressSpace() != LangAS::Default) { - if (Type.getAddressSpace() != ASIdx) { - S.Diag(Attr.getLoc(), diag::err_attribute_address_multiple_qualifiers); - Attr.setInvalid(); - return; - } else - // Emit a warning if they are identical; it's likely unintended. - S.Diag(Attr.getLoc(), - diag::warn_attribute_address_multiple_identical_qualifiers); + if (DiagnoseMultipleAddrSpaceAttributes(S, Type.getAddressSpace(), ASIdx, + Attr.getLoc())) { + Attr.setInvalid(); + return; } Type = S.Context.getAddrSpaceQualType(Type, ASIdx); Index: test/CodeGenOpenCLCXX/method-overload-address-space.cl =================================================================== --- /dev/null +++ test/CodeGenOpenCLCXX/method-overload-address-space.cl @@ -0,0 +1,35 @@ +//RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -O0 -o - | FileCheck %s + +struct C { + void foo() __local; + void foo() __global; + void foo(); + void bar(); +}; + +__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)* + c1.foo(); + // CHECK: call void @_ZNU3AS31C3fooEv(%struct.C addrspace(3)* + c2.foo(); + // CHECK: call void @_ZNU3AS41C3fooEv(%struct.C addrspace(4)* + c3.foo(); + // CHECK: call void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)* + c_ptr->foo(); + // CHECK: 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)*)) + 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)*)) + c_ref.bar(); +} Index: test/SemaOpenCLCXX/method-overload-address-space.cl =================================================================== --- /dev/null +++ test/SemaOpenCLCXX/method-overload-address-space.cl @@ -0,0 +1,20 @@ +//RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -pedantic -verify + +struct C { + void m1() __local __local; //expected-warning{{multiple identical address spaces specified for type}} + //expected-note@-1{{candidate function}} + void m1() __global; + //expected-note@-1{{candidate function}} + void m2() __global __local; //expected-error{{multiple address spaces specified for type}} +}; + +__global C c_glob; + +__kernel void bar() { + __local C c_loc; + C c_priv; + + c_glob.m1(); + c_loc.m1(); + c_priv.m1(); //expected-error{{no matching member function for call to 'm1'}} +}