Index: lib/Parse/ParseDecl.cpp =================================================================== --- lib/Parse/ParseDecl.cpp +++ lib/Parse/ParseDecl.cpp @@ -6143,6 +6143,24 @@ RestrictQualifierLoc = DS.getRestrictSpecLoc(); } + // Propagate address space attributes with a language semantic. + if (getLangOpts().OpenCLCPlusPlus) { + for (ParsedAttr &attr : DS.getAttributes()) { + switch (attr.getKind()) { + case ParsedAttr::AT_OpenCLConstantAddressSpace: + case ParsedAttr::AT_OpenCLLocalAddressSpace: + case ParsedAttr::AT_OpenCLGlobalAddressSpace: + case ParsedAttr::AT_OpenCLPrivateAddressSpace: + case ParsedAttr::AT_OpenCLGenericAddressSpace: + FnAttrs.addNew(attr.getName(), attr.getLoc(), nullptr, + attr.getLoc(), nullptr, 0, ParsedAttr::AS_Keyword); + break; + default: + break; + } + } + } + // Parse ref-qualifier[opt]. if (ParseRefQualifier(RefQualifierIsLValueRef, RefQualifierLoc)) EndLoc = RefQualifierLoc; Index: lib/Sema/SemaOverload.cpp =================================================================== --- lib/Sema/SemaOverload.cpp +++ lib/Sema/SemaOverload.cpp @@ -1142,18 +1142,19 @@ // 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(); if (!getLangOpts().CPlusPlus14 && NewMethod->isConstexpr() && !isa(NewMethod)) NewQuals |= Qualifiers::Const; - // We do not allow overloading based off of '__restrict'. OldQuals &= ~Qualifiers::Restrict; NewQuals &= ~Qualifiers::Restrict; if (OldQuals != NewQuals) return true; + if (OldMethod->getTypeQualifiers().getAddressSpace() != + NewMethod->getTypeQualifiers().getAddressSpace()) + return true; } // Though pass_object_size is placed on parameters and takes an argument, we @@ -6654,6 +6655,20 @@ Candidate.Viable = false; Candidate.FailureKind = ovl_non_default_multiversion_function; } + + if (Method->getType().getQualifiers().hasAddressSpace() && + !ObjectType.isNull()) { + Qualifiers QualsObject = + ObjectType->isPointerType() || ObjectType->isReferenceType() + ? ObjectType->getPointeeType().getQualifiers() + : ObjectType.getQualifiers(); + Qualifiers QualsMethod = Method->getType().getQualifiers(); + if (!QualsMethod.isAddressSpaceSupersetOf(QualsObject)) { + Candidate.Viable = false; + Candidate.FailureKind = ovl_fail_bad_final_conversion; + return; + } + } } /// Add a C++ member function template as a candidate to the candidate @@ -9256,6 +9271,14 @@ if (HasPS1 != HasPS2 && HasPS1) return true; + if (S.getLangOpts().OpenCLCPlusPlus && Cand1.Function && Cand2.Function) { + LangAS CandAS1 = Cand1.Function->getType().getAddressSpace(); + LangAS CandAS2 = Cand2.Function->getType().getAddressSpace(); + if ((CandAS2 == LangAS::opencl_generic || CandAS2 == LangAS::Default) && + (CandAS1 != LangAS::opencl_generic && CandAS1 != LangAS::Default)) + return true; + } + return isBetterMultiversionCandidate(Cand1, Cand2); } Index: lib/Sema/SemaType.cpp =================================================================== --- lib/Sema/SemaType.cpp +++ lib/Sema/SemaType.cpp @@ -3921,6 +3921,14 @@ llvm_unreachable("unknown NullabilityKind"); } +/// IsClassMember - Determines whether a state belongs to a class member. +static bool IsClassMember(TypeProcessingState &State) { + return (!State.getDeclarator().getCXXScopeSpec().isEmpty() && + State.getDeclarator().getCXXScopeSpec().getScopeRep()->getKind() == + NestedNameSpecifier::TypeSpec) || + State.getDeclarator().getContext() == DeclaratorContext::MemberContext; +} + static TypeSourceInfo * GetTypeSourceInfoForDeclarator(TypeProcessingState &State, QualType T, TypeSourceInfo *ReturnTypeInfo); @@ -4825,18 +4833,45 @@ 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(); + IsClassMember(state)) { + LangAS ASIdx = LangAS::Default; + // Take address space attr if any and mark as invalid to avoid adding + // them later while creating QualType. + for (ParsedAttr &attr : DeclType.getAttrs()) { + switch (attr.getKind()) { + case ParsedAttr::AT_OpenCLConstantAddressSpace: + ASIdx = LangAS::opencl_constant; + attr.setInvalid(); + break; + case ParsedAttr::AT_OpenCLLocalAddressSpace: + ASIdx = LangAS::opencl_local; + attr.setInvalid(); + break; + case ParsedAttr::AT_OpenCLGlobalAddressSpace: + ASIdx = LangAS::opencl_global; + attr.setInvalid(); + break; + case ParsedAttr::AT_OpenCLPrivateAddressSpace: + ASIdx = LangAS::opencl_private; + attr.setInvalid(); + break; + case ParsedAttr::AT_OpenCLGenericAddressSpace: + ASIdx = LangAS::opencl_generic; + attr.setInvalid(); + break; + default: + break; + } + if (ASIdx != LangAS::Default) + break; + } // 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); T = state.getSema().Context.getAddrSpaceQualType(T, AS); @@ -5834,7 +5869,10 @@ // ISO/IEC TR 18037 S5.3 (amending C99 6.7.3): "A function type shall not be // qualified by an address-space qualifier." - if (Type->isFunctionType()) { + // Allow qualifying methods as an extension. + // FIXME: For now only enabled for OpenCL. + if (Type->isFunctionType() && + !(S.getLangOpts().OpenCLCPlusPlus && IsClassMember(State))) { S.Diag(Attr.getLoc(), diag::err_attribute_address_function_type); Attr.setInvalid(); return; 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'}} +}