Index: include/clang/AST/Type.h =================================================================== --- include/clang/AST/Type.h +++ include/clang/AST/Type.h @@ -1982,7 +1982,7 @@ bool isObjCQualifiedClassType() const; // Class bool isObjCObjectOrInterfaceType() const; bool isObjCIdType() const; // id - + bool isDecltypeType() const; /// Was this type written with the special inert-in-ARC __unsafe_unretained /// qualifier? /// @@ -6440,6 +6440,10 @@ return isObjCIdType() || isObjCClassType() || isObjCSelType(); } +inline bool Type::isDecltypeType() const { + return isa(this); +} + #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ inline bool Type::is##Id##Type() const { \ return isSpecificBuiltinType(BuiltinType::Id); \ Index: include/clang/Sema/ParsedAttr.h =================================================================== --- include/clang/Sema/ParsedAttr.h +++ include/clang/Sema/ParsedAttr.h @@ -567,6 +567,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 @@ -6177,6 +6177,20 @@ Qualifiers Q = Qualifiers::fromCVRUMask(DS.getTypeQualifiers()); if (D.getDeclSpec().isConstexprSpecified() && !getLangOpts().CPlusPlus14) Q.addConst(); + // FIXME: Collect C++ address spaces. + // If there are multiple different address spaces, the source is invalid. + // Carry on using the first addr space for the qualifiers of 'this'. + // The diagnostic will be given later while creating the function + // prototype for the method. + 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 @@ -1171,16 +1171,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; } @@ -5150,6 +5148,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 @@ -3915,6 +3915,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; + } + // 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); @@ -4822,18 +4841,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); @@ -5789,19 +5825,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); } @@ -5879,34 +5905,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); @@ -7243,9 +7249,12 @@ // Do not deduce addr space of the void type, e.g. in f(void), otherwise // it will fail some sema check. (T->isVoidType() && !IsPointee) || - // Do not deduce address spaces for dependent types because they might end + // Do not deduce addr spaces for dependent types because they might end // up instantiating to a type with an explicit address space qualifier. - T->isDependentType()) + T->isDependentType() || + // Do not deduce addr space of decltype because it will be taken from + // its argument. + T->isDecltypeType()) return; LangAS ImpAddr = LangAS::Default; Index: test/CodeGenOpenCLCXX/method-overload-address-space.cl =================================================================== --- test/CodeGenOpenCLCXX/method-overload-address-space.cl +++ 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/address-space-of-this-class-scope.cl =================================================================== --- test/SemaOpenCLCXX/address-space-of-this-class-scope.cl +++ test/SemaOpenCLCXX/address-space-of-this-class-scope.cl @@ -0,0 +1,18 @@ +//RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -pedantic -verify + +struct C { + auto fGlob() __global -> decltype(this); + auto fGen() -> decltype(this); + auto fErr() __global __local -> decltype(this); //expected-error{{multiple address spaces specified for type}} +}; + +void bar(__local C*); +// expected-note@-1{{candidate function not viable: address space mismatch in 1st argument ('decltype(this)' (aka '__global C *')), parameter type must be '__local C *'}} +// expected-note@-2{{candidate function not viable: address space mismatch in 1st argument ('decltype(this)' (aka 'C *')), parameter type must be '__local C *'}} + +__global C Glob; +void foo(){ +bar(Glob.fGlob()); // expected-error{{no matching function for call to 'bar'}} +// FIXME: AS of 'this' below should be correctly deduced to generic +bar(Glob.fGen()); // expected-error{{no matching function for call to 'bar'}} +} Index: test/SemaOpenCLCXX/address_space_overloading.cl =================================================================== --- test/SemaOpenCLCXX/address_space_overloading.cl +++ test/SemaOpenCLCXX/address_space_overloading.cl @@ -1,12 +1,12 @@ // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=c++ -// expected-no-diagnostics +// FIXME: This test shouldn't trigger any errors. struct RetGlob { int dummy; }; -struct RetGen { +struct RetGen { //expected-error{{binding value of type '__generic RetGen' to reference to type 'RetGen' drops <> qualifiers}} char dummy; }; @@ -19,5 +19,5 @@ __local int *ArgLoc; RetGlob TestGlob = foo(ArgGlob); RetGen TestGen = foo(ArgGen); - TestGen = foo(ArgLoc); + TestGen = foo(ArgLoc); //expected-note{{in implicit copy assignment operator for 'RetGen' first required here}} } Index: test/SemaOpenCLCXX/method-overload-address-space.cl =================================================================== --- test/SemaOpenCLCXX/method-overload-address-space.cl +++ 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'}} +}