[OpenCL] Allow address spaces as method qualifiers

[Anastasia]

Summary

This patch allows qualifying methods with address spaces and extends some overloading rules to use those qualifiers.

The main use case is to prevent conversions to generic/default address space where it's requested by the programmer. More details can be found in:
http://lists.llvm.org/pipermail/cfe-dev/2018-December/060470.html

This patch doesn't enable the feature for C++ yet but it can easily be generalized if the overall flow is right.

[Anastasia]

Removed FIXME that has been addressed

[Anastasia]

lib/Sema/SemaOverload.cpp

2828

I am still missing something here.

[rjmccall]

You're gating on OpenCL C++ in several places in this patch, but I don't think you need to. This extension should be available to anyone using address spaces and C++.

lib/Parse/ParseDecl.cpp

6162

This is enforcing a restriction that users write const __private, which seems unreasonable. It looks like ParseTypeQualifierList takes a flag saying not to parse attributes; try adding a new option to that enum allowing address-space attributes.

Collecting the attributes on the type-qualifiers DeclSpec rather than adding them as function attributes seems correct.

lib/Sema/SemaOverload.cpp

1157

If you leave OldQuals and NewQuals as Qualifiers and then just remove restrict from both sets, I think you don't need a special check for address spaces here.

6671

I would expect this to fail in TryObjectArgumentInitialization above and not to need a special case here.

9279

This at least needs a comment explaining the rule you're trying to implement.

[Anastasia]

Address review comments.

lib/Parse/ParseDecl.cpp

6162

Do you mean ParseTypeQualifierListOpt? That already parses the address spaces unconditionally. The problem is however that we are using local DeclSpec - DS variable here during the function qualifiers parsing because the DeclSpec member of the Declarator corresponds to the return type as far as I understand it. Therefore I am propagating missing address space attributes from the local DS variable into FnAttrs to be used in the function type info.

With current patch both of the following two forms:

struct C {
  void foo() __local const;
};

and

struct C {
  void foo() const __local;
};

would be parsed correctly generating identical IR

declare void @_ZNU3AS3K1C3fooEv(%struct.C addrspace(3)*)

[Anastasia]

In D55850#1335826, @rjmccall wrote:

You're gating on OpenCL C++ in several places in this patch, but I don't think you need to. This extension should be available to anyone using address spaces and C++.

Ok, I am happy to generalize it to C++. There are some slight differences in the address space attribute representation for C/C++ as I am now figuring out. So I would prefer to put a separate patch if it's ok with you.

[rjmccall]

lib/Parse/ParseDecl.cpp

6162

Oh, I see, sorry. Why filter on attributes at all, then? We should *only* be parsing qualifier attributes in that list.

I actually think it would be reasonable to change DeclaratorChunk::FunctionTypeInfo to just store a DeclSpec for all the qualifiers; we're already duplicating an unfortunate amount of the logic from DeclSpec, like remembering SourceLocations for all the qualifiers. You'll have to store it out-of-line, but we already store e.g. parameters out of line, so the machinery for allocating and destroying it already exists. Just make sure we don't actually allocate anything in the common case where there aren't any qualifiers (i.e. for C and non-member C++ functions).

Also, I suspect that the use of CXXThisScopeRAII below this needs to be updated to pull *all* the qualifiers out of DS. Maybe Sema should have a function for that.

lib/Sema/SemaOverload.cpp

2828

Well, at least the failure here is just to fall into the generic diagnostic.

Getting this diagnostic right probably requires some minor work to the diagnostics engine. If you look at err_init_conversion_failed, which is (I think) the diagnostic that's always being used here, it matches every possible CVR mask so that it can pretty-print them. This is already a problem because the input is actually a CVRU mask! A better option would be to teach DiagnosticEngine how to store and format a Qualifiers value, and then you can just stream the original Qualifiers into the diagnostic here.

But that's obviously work for a separate patch.

9279

Okay, thanks for the clarification. I think this should just be part of CompareImplicitConversionSequence, right? It seems to me that you should prefer e.g. int __private * -> int __private * over int __generic * as a normal argument conversion as well.

Also, can this be written in terms of isAddressSpaceSupersetOf? I don't remember how LangAS::Default works in OpenCL C++ enough to understand how it fits in here.

[Anastasia]

lib/Parse/ParseDecl.cpp

6162

I have uploaded a separate patch for this:
https://reviews.llvm.org/D55948

I think I can't avoid storing DeclSpec for non-memeber functions in C++ because we still use them to give diagnostics about the qualifiers. For example:

test.cpp:1:12: error: non-member function cannot have 'const' qualifier
void foo() const;

As for CXXThisScopeRAII, we seem to be adding other qualifiers to the QualType directly (while building it), so there seems to be no function to collect them into Qualifiers. I can, however, add code to collect address space qualifiers here. I guess we don't have any other missing qualifiers to collect?

[rjmccall]

lib/Parse/ParseDecl.cpp

6162

I think you can probably avoid allocating and storing a DeclSpec if there are no qualifiers, which should be easy to test, and then the absence of a DeclSpec should allow Sema to quickly bail out of the check for illegal qualifiers on a non-member function type.

I guess we don't have any other missing qualifiers to collect?

Not that are legal to apply to this.

[Anastasia]

lib/Sema/SemaOverload.cpp

9279

That's correct we should implement the same logic for the arguments too. I will create a helper function or do you think we should just call CompareImplicitConversionSequence on the method type too?

I think isAddressSpaceSupersetOf can't be used here because it determines compatibility of address spaces. However, the logic we are expressing is for the address space preference instead.

[rjmccall]

lib/Sema/SemaOverload.cpp

9279

If I understand correctly, we already call CompareImplicitConversionSequence on the object-argument conversion above, as part of the for (unsigned ArgIdx = StartArg; ArgIdx < NumArgs; ++ArgIdx) loop.

I think the address-space ordering rule can be reasonably based on compatibility. In fact, I think it already is in our implementation. The right rule is basically that (1) an exact match is better than a conversion and (2) a conversion to a subspace is better than a conversion to a superspace. You can think of this as modifying the "proper subset" rule of [over.ics.rank]p3.2.5, which in implementation terms means Qualifiers::isMoreQualifiedThan, which already ends up using isAddressSpaceSupersetOf. So if this ranking isn't working right, it's probably because we're incorrectly fast-pathing based on CVR qualifiers somewhere, and in fact I can see a pretty suspicious check like that in CompareQualificationConversions.

[Anastasia]

lib/Sema/SemaOverload.cpp

2828

I created a patch for this in https://reviews.llvm.org/D56198.

[brunodf]

lib/Sema/SemaType.cpp

4877

I follow all of the above (from the point "a class member function has an address space"), but then I take issue with this line (already from Mikael).

You look at the declaration's attributes, to derive the ASIdx relating to the method's this argument. You mark the relevant attributes as invalid, to prevent them from being considered in "processTypeAttrs" after the switch that we break below. The ASIdx is stored in the qualifiers EPI to go to the FunctionProtoType (this will affect getThisType). This all seems fine.

But then this line adds the address space qualification to T, the type of the declared function itself. This seems unnecessary, and conceptually wrong: while the function may have an address space in syntax, this address space applies to the this argument, not to the function object (and a pointer to the function is not a pointer to this address space etc.). In short, I think this line should be removed.

[Anastasia]

Rebased on top of recent related changes and addressed remaining review comments.

This now depends on: https://reviews.llvm.org/D56735

[Anastasia]

lib/Sema/SemaOverload.cpp

9279

I created a separate review to fix this: https://reviews.llvm.org/D56735

lib/Sema/SemaType.cpp

4877

I think I fixed it in: https://reviews.llvm.org/D56066

[rjmccall]

Okay, so is this ready to re-review independently?

[Anastasia]

In D55850#1358812, @rjmccall wrote:

Okay, so is this ready to re-review independently?

Yes, please. It should be fine to review on its own. Thanks!

[Anastasia]

Fixed wording on the comment.

[rjmccall]

lib/Parse/ParseDecl.cpp

6195

Does this not need to diagnose redundant qualifiers? Why is this path required in addition to the path in SemaType, anyway?

lib/Sema/SemaType.cpp

3935

I'd just drop this else, but if you want to include it, please use braces for consistency with the rest of the if-chain.

[Anastasia]

• Removed else case in DiagnoseMultipleAddrSpaceAttributes
• Added extra test to check correctness of addr space for 'this' when statically used in a class scope.

[Anastasia]

lib/Parse/ParseDecl.cpp

6195

We discussed earlier to collect addr space qualifiers for completeness: https://reviews.llvm.org/D55850#inline-495037

Without this change the following doesn't work for addr spaces correctly:

auto fGlob() __global -> decltype(this);

I added a test that check this now in test/SemaOpenCLCXX/address-space-of-this-class-scope.cl.

Here we only collect the quals to be applied to 'this'. If there are multiple address space specified on a method it will be diagnosed in SemaType.cpp. I think we probably don't need to give the same diagnostic twice?

[Anastasia]

• Fixed typo in FIXME

[rjmccall]

lib/Parse/ParseDecl.cpp

6195

Okay, if this is already diagnosed, it isn't a problem. Please leave a comment to that effect.

[Anastasia]

Added a comment explaining when multiple address spaces are diagnosed.

[rjmccall]

Okay, LGTM.

[Anastasia]

After rebase I had to modify the following test:

Index: test/SemaOpenCLCXX/address_space_overloading.cl
===================================================================
--- test/SemaOpenCLCXX/address_space_overloading.cl	(revision 351746)
+++ test/SemaOpenCLCXX/address_space_overloading.cl	(working copy)
@@ -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 <<ERROR>> 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}}
 }

After looking at it, I realized that there is another un-handled path for address spaces in the initialization sequence. That gets hit during the creation of a return statement when we define implicit copy assignment. assignment. As a result I end up with the following incorrect AST:

UnaryOperator 0x75a910 '__generic struct RetGen' lvalue prefix '*' cannot overflow
`-CXXThisExpr 0x75a900 '__generic struct RetGen *' this

As I wasn't sure whether I should fix the initialization sequence again by splitting the address space conversion to move it to a later step or just fix the type of *this expr, I will upload a separate fix for this instead of reopening this review.

The overloading resolution seems to work fine at least which is what this test is supposed to check.

[ebevhan]

I'm a bit late to the party here, it was an older patch so I wasn't watching this one.

I get a bit miffed when address space related features get locked behind langoptions that aren't strictly address spaces. I know that there are currently a couple code snippets which are behind LangOptions.OpenCL, that are needed for address spaces to work reasonably even when you aren't using OpenCL.

I guess I do understand that the only address spaces that are interesting to parse here are the named qualifier ones, but it would be convenient if the parsing would accept other ones as well (such as the __attribute__ based ones) and not necessarily assume that the user is using OpenCL++.

EDIT: Sort of jumped the gun there... The summary even says "This patch doesn't enable the feature for C++ yet but it can easily be generalized if the overall flow is right." Don't mind me.

[Anastasia]

In D55850#1365437, @ebevhan wrote:

I'm a bit late to the party here, it was an older patch so I wasn't watching this one.

I get a bit miffed when address space related features get locked behind langoptions that aren't strictly address spaces. I know that there are currently a couple code snippets which are behind LangOptions.OpenCL, that are needed for address spaces to work reasonably even when you aren't using OpenCL.

I guess I do understand that the only address spaces that are interesting to parse here are the named qualifier ones, but it would be convenient if the parsing would accept other ones as well (such as the __attribute__ based ones) and not necessarily assume that the user is using OpenCL++.

EDIT: Sort of jumped the gun there... The summary even says "This patch doesn't enable the feature for C++ yet but it can easily be generalized if the overall flow is right." Don't mind me.

This is correct. My plan was to generalize to C++ in the next patch. As I am more familiar with OpenCL and many rules are already well defined, this was my starting point. However, I want this work to be used for C++ based implementations generically and not just for OpenCL. I am still struggling to understand the semantic of address spaces in C++ for some cases and need a little bit more time, but I think I should be able to upload the initial patch soon. I would then appreciate any help with reviewing to get it right!

[ebevhan]

Okay, sounds good! I'm not a C++ expert, but I'd be happy to look at the patches when they're up. Haven't done much serious testing on my end so far, but from what I've seen the patches so far look good!

[Anastasia]

In D55850#1366094, @ebevhan wrote:

Okay, sounds good! I'm not a C++ expert, but I'd be happy to look at the patches when they're up. Haven't done much serious testing on my end so far, but from what I've seen the patches so far look good!

Cool! Thanks a lot for your feedback btw. Let me know if you discover any issues/bugs! I do need to improve testing at some point soon as I am sure there are still a lot of uncaught corner cases.

[rjmccall]

I'll just add that we've been trying to not put things behind OpenCL guards as much as possible. Most of the remaining OpenCL checks are for OpenCL-specific logic like inferring the default address space, and yeah, we could probably make that a target option or something.

[Anastasia]

In D55850#1366709, @rjmccall wrote:

Most of the remaining OpenCL checks are for OpenCL-specific logic like inferring the default address space, and yeah, we could probably make that a target option or something.

We have an address space map in the targets that can be used to map default address space if needed. Or are you suggesting to migrate functionality like deduceOpenCLImplicitAddrSpace() to the target setting? Although I feel this belong to the language rather.

[rjmccall]

In D55850#1368767, @Anastasia wrote:

In D55850#1366709, @rjmccall wrote:

Most of the remaining OpenCL checks are for OpenCL-specific logic like inferring the default address space, and yeah, we could probably make that a target option or something.

We have an address space map in the targets that can be used to map default address space if needed. Or are you suggesting to migrate functionality like deduceOpenCLImplicitAddrSpace() to the target setting? Although I feel this belong to the language rather.

True. But we could make it a Sema-wide setting, configured immediately so that we don't have to repeatedly check a language option whenever we parse or synthesize a method. It would make that logic feel less language-specific.

[ebevhan]

I think the code that comes to mind is mostly like in GetTypeSourceInfoForDeclarator:

LangAS AS =
    (ASIdx == LangAS::Default ? LangAS::opencl_generic : ASIdx);

It's behind an OpenCLCPlusPlus guard so it won't add generic on anything if it's not OpenCL++, but there shouldn't be a reason why the rest of the code in that block won't work for regular C++.

In fact, in most regular C++ this would be an issue, since Default typically _is_ the 'generic' address space there.

[ebevhan]

lib/Sema/SemaOverload.cpp

5124

I tested this patch with some kludges to let it parse our address space qualifiers, and hit a problem here; just because the FromType isn't qualified doesn't mean that the object conversion is okay. A conversion of an object Type to __X Type might not be legal if there is no conversion from 'no address space' to '__X' address space.

The example was a class with an AS-qualified overload, and when resolving for calling the overloaded method on an object Type *, it hit an ambiguous resolution since it considered the AS-qualified method to be a legal candidate.

It feels like this might be rather specific to how a language/target wants to use address spaces, though.

[Anastasia]

In D55850#1369008, @ebevhan wrote:

I think the code that comes to mind is mostly like in GetTypeSourceInfoForDeclarator:

LangAS AS =
    (ASIdx == LangAS::Default ? LangAS::opencl_generic : ASIdx);

It's behind an OpenCLCPlusPlus guard so it won't add generic on anything if it's not OpenCL++, but there shouldn't be a reason why the rest of the code in that block won't work for regular C++.

In fact, in most regular C++ this would be an issue, since Default typically _is_ the 'generic' address space there.

Yes, opencl_generic is something quite specific to OpenCL I feel. I am not sure we can generalize it to C++. I would quite like to understand this better. However, I think we could potentially implement opencl_generic using Default if only earlier OpenCL implementations wouldn't use Default for opencl_private. And because everything was built on top of this assumption it's now a big and a scary refactoring change to make. However, I think we can align different language modes better if opencl_generic could be implemented as Default. I might look into this at some point.

lib/Parse/ParseDecl.cpp

if (getLangOpts().CPlusPlus) {
                                 }));
       if (!DS.getSourceRange().getEnd().isInvalid()) {
         EndLoc = DS.getSourceRange().getEnd();
         ConstQualifierLoc = DS.getConstSpecLoc();
         VolatileQualifierLoc = DS.getVolatileSpecLoc();
         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;
 
       // C++11 [expr.prim.general]p3:
       //   If a declaration declares a member function or member function
       //   template of a class X, the expression this is a prvalue of type
       //   "pointer to cv-qualifier-seq X" between the optional cv-qualifer-seq
if (getLangOpts().CPlusPlus) {
 
       Sema::CXXThisScopeRAII ThisScope(
           Actions, dyn_cast<CXXRecordDecl>(Actions.CurContext), Q,
           IsCXX11MemberFunction);
 
       // Parse exception-specification[opt].
       bool Delayed = D.isFirstDeclarationOfMember() &&
                      D.isFunctionDeclaratorAFunctionDeclaration();
       if (Delayed && Actions.isLibstdcxxEagerExceptionSpecHack(D) &&
           GetLookAheadToken(0).is(tok::kw_noexcept) &&
           GetLookAheadToken(1).is(tok::l_paren) &&
           GetLookAheadToken(2).is(tok::kw_noexcept) &&
           GetLookAheadToken(3).is(tok::l_paren) &&
           GetLookAheadToken(4).is(tok::identifier) &&
           GetLookAheadToken(4).getIdentifierInfo()->isStr("swap")) {
         // HACK: We've got an exception-specification
         //   noexcept(noexcept(swap(...)))

lib/Sema/SemaOverload.cpp

if (OldMethod->getRefQualifier() != NewMethod->getRefQualifier()) {
       }
       return true;
     }
 
     // We may not have applied the implicit const for a constexpr member
     // 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<CXXConstructorDecl>(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
   // consider it to be a function-level modifier for the sake of function
   // identity. Either the function has one or more parameters with
   // pass_object_size or it doesn't.
   if (functionHasPassObjectSizeParams(New) !=
       functionHasPassObjectSizeParams(Old))
void Sema::HandleFunctionTypeMismatch(PartialDiagnostic &PDiag,
   // Handle different return type.
   if (!Context.hasSameType(FromFunction->getReturnType(),
                            ToFunction->getReturnType())) {
     PDiag << ft_return_type << ToFunction->getReturnType()
           << FromFunction->getReturnType();
     return;
   }
 
   // FIXME: OpenCL: Need to consider address spaces
   unsigned FromQuals = FromFunction->getTypeQuals().getCVRUQualifiers();
   unsigned ToQuals = ToFunction->getTypeQuals().getCVRUQualifiers();
   if (FromQuals != ToQuals) {
     PDiag << ft_qualifer_mismatch << ToQuals << FromQuals;
     return;
   }
 
   // Handle exception specification differences on canonical type (in C++17
TryObjectArgumentInitialization(Sema &S, SourceLocation Loc, QualType FromType,
   if (ImplicitParamType.getCVRQualifiers()
                                     != FromTypeCanon.getLocalCVRQualifiers() &&
       !ImplicitParamType.isAtLeastAsQualifiedAs(FromTypeCanon)) {
     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);
   ImplicitConversionKind SecondKind;
   if (ClassTypeCanon == FromTypeCanon.getLocalUnqualifiedType()) {
     SecondKind = ICK_Identity;
   } else if (S.IsDerivedFrom(Loc, FromType, ClassType))
     SecondKind = ICK_Derived_To_Base;
   else {
if (EnableIfAttr *FailedAttr = CheckEnableIf(Method, Args, true)) {
     return;
   }
 
   if (Method->isMultiVersion() && Method->hasAttr<TargetAttr>() &&
       !Method->getAttr<TargetAttr>()->isDefaultVersion()) {
     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
 /// set, using template argument deduction to produce an appropriate member
 /// function template specialization.
 void
 Sema::AddMethodTemplateCandidate(FunctionTemplateDecl *MethodTmpl,
                                  DeclAccessPair FoundDecl,
bool clang::isBetterOverloadCandidate(
 
   bool HasPS1 = Cand1.Function != nullptr &&
                 functionHasPassObjectSizeParams(Cand1.Function);
   bool HasPS2 = Cand2.Function != nullptr &&
                 functionHasPassObjectSizeParams(Cand2.Function);
   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);
 }
 
 /// Determine whether two declarations are "equivalent" for the purposes of
 /// name lookup and overload resolution. This applies when the same internal/no
 /// linkage entity is defined by two modules (probably by textually including
 /// the same header). In such a case, we don't consider the declarations to
 /// declare the same entity, but we also don't want lookups with both

lib/Sema/SemaType.cpp

case NullabilityKind::Nullable:
     return createSimpleAttr<TypeNullableAttr>(Ctx, Attr);
 
   case NullabilityKind::Unspecified:
     return createSimpleAttr<TypeNullUnspecifiedAttr>(Ctx, Attr);
   }
   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);
 
 static TypeSourceInfo *GetFullTypeForDeclarator(TypeProcessingState &state,
                                                 QualType declSpecType,
                                                 TypeSourceInfo *TInfo) {
   // The TypeSourceInfo that this function returns will not be a null type.
   // If there is an error, this function will fill in a dummy type as fallback.
   QualType T = declSpecType;
   Declarator &D = state.getDeclarator();
   Sema &S = state.getSema();
case DeclaratorChunk::Function: {
         S.checkExceptionSpecification(D.isFunctionDeclarationContext(),
                                       FTI.getExceptionSpecType(),
                                       DynamicExceptions,
                                       DynamicExceptionRanges,
                                       NoexceptExpr,
                                       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);
         } else {
           T = Context.getFunctionType(T, ParamTys, EPI);
         }
       }
       break;
     }
     case DeclaratorChunk::MemberPointer: {
       // The scope spec must refer to a class, or be dependent.

 /// space for the type.
 static void HandleAddressSpaceTypeAttribute(QualType &Type,
                                             const ParsedAttr &Attr,
                                             TypeProcessingState &State) {
   Sema &S = State.getSema();
 
   // 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;
   }
 
   LangAS ASIdx;
   if (Attr.getKind() == ParsedAttr::AT_AddressSpace) {
 

test/CodeGenOpenCLCXX/method-overload-address-space.cl

+//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();
+}

test/SemaOpenCLCXX/method-overload-address-space.cl

+//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'}}
+}