[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

2853–2854

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

6173

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

1186

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.

6708

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

9311

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

[Anastasia]

Address review comments.

lib/Parse/ParseDecl.cpp

6173

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

6173

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

2853–2854

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.

9311

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

6173

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

6173

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

9311

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

9311

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

2853–2854

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

[brunodf]

lib/Sema/SemaType.cpp

4873–4874

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

9311

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

lib/Sema/SemaType.cpp

4873–4874

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

6188

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

3929

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

6188

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

6188

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

5151

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.

include/clang/AST/Type.h

public:
   // FIXME: change this to 'raw' interface type, so we can used 'interface' type
   // for the common case.
   bool isObjCObjectType() const;                // NSString or typeof(*(id)0)
   bool isObjCQualifiedInterfaceType() const;    // NSString<foo>
   bool isObjCQualifiedIdType() const;           // id<foo>
   bool isObjCQualifiedClassType() const;        // Class<foo>
   bool isObjCObjectOrInterfaceType() const;
   bool isObjCIdType() const;                    // id
-
+  bool isDecltypeType() const;
   /// Was this type written with the special inert-in-ARC __unsafe_unretained
   /// qualifier?
   ///
   /// This approximates the answer to the following question: if this
   /// translation unit were compiled in ARC, would this type be qualified
   /// with __unsafe_unretained?
   bool isObjCInertUnsafeUnretainedType() const {
     return hasAttr(attr::ObjCInertUnsafeUnretained);
if (const auto *OPT = getAs<PointerType>())
     return OPT->getPointeeType()->isSpecificBuiltinType(BuiltinType::ObjCSel);
   return false;
 }
 
 inline bool Type::isObjCBuiltinType() const {
   return isObjCIdType() || isObjCClassType() || isObjCSelType();
 }
 
+inline bool Type::isDecltypeType() const {
+  return isa<DecltypeType>(this);
+}
+
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   inline bool Type::is##Id##Type() const { \
     return isSpecificBuiltinType(BuiltinType::Id); \
   }
 #include "clang/Basic/OpenCLImageTypes.def"
 
 inline bool Type::isSamplerT() const {
   return isSpecificBuiltinType(BuiltinType::OCLSampler);

include/clang/Sema/ParsedAttr.h

public:
   bool isSupportedByPragmaAttribute() const;
 
   /// If the parsed attribute has a semantic equivalent, and it would
   /// have a semantic Spelling enumeration (due to having semantically-distinct
   /// spelling variations), return the value of that semantic spelling. If the
   /// 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;
 /// A factory, from which one makes pools, from which one creates
 /// individual attributes which are deallocated with the pool.
 ///
 /// Note that it's tolerably cheap to create and destroy one of
 /// these as long as you don't actually allocate anything in it.

lib/Parse/ParseDecl.cpp

if (getLangOpts().CPlusPlus) {
       //   and the end of the function-definition, member-declarator, or
       //   declarator.
       // FIXME: currently, "static" case isn't handled correctly.
       bool IsCXX11MemberFunction =
         getLangOpts().CPlusPlus11 &&
         D.getDeclSpec().getStorageClassSpec() != DeclSpec::SCS_typedef &&
         (D.getContext() == DeclaratorContext::MemberContext
          ? !D.getDeclSpec().isFriendSpecified()
          : D.getContext() == DeclaratorContext::FileContext &&
            D.getCXXScopeSpec().isValid() &&
            Actions.CurContext->isRecord());
 
       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<CXXRecordDecl>(Actions.CurContext), Q,
           IsCXX11MemberFunction);
 
       // Parse exception-specification[opt].
       bool Delayed = D.isFirstDeclarationOfMember() &&
                      D.isFunctionDeclaratorAFunctionDeclaration();

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();
+    auto OldQuals = OldMethod->getTypeQualifiers();
+    auto NewQuals = NewMethod->getTypeQualifiers();
     if (!getLangOpts().CPlusPlus14 && NewMethod->isConstexpr() &&
         !isa<CXXConstructorDecl>(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;
   }
 
   // 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))
     return true;
 
   // enable_if attributes are an order-sensitive part of the signature.
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;
   }
 
   if (FromFunction->getTypeQuals() != ToFunction->getTypeQuals()) {
     PDiag << ft_qualifer_mismatch << ToFunction->getTypeQuals()
           << FromFunction->getTypeQuals();
     return;
   }
 
   // Handle exception specification differences on canonical type (in C++17
   // onwards).
   if (cast<FunctionProtoType>(FromFunction->getCanonicalTypeUnqualified())
TryObjectArgumentInitialization(Sema &S, SourceLocation Loc, QualType FromType,
   if (ImplicitParamType.getCVRQualifiers()
                                     != FromTypeCanon.getLocalCVRQualifiers() &&
       !ImplicitParamType.isAtLeastAsQualifiedAs(FromTypeCanon)) {
     ICS.setBad(BadConversionSequence::bad_qualifiers,
                FromType, ImplicitParamType);
     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);
   ImplicitConversionKind SecondKind;
   if (ClassTypeCanon == FromTypeCanon.getLocalUnqualifiedType()) {
     SecondKind = ICK_Identity;
   } else if (S.IsDerivedFrom(Loc, FromType, ClassType))
     SecondKind = ICK_Derived_To_Base;

 void
 Sema::AddMethodTemplateCandidate(FunctionTemplateDecl *MethodTmpl,
                                  DeclAccessPair FoundDecl,
                                  CXXRecordDecl *ActingContext,
                                  TemplateArgumentListInfo *ExplicitTemplateArgs,
                                  QualType ObjectType,
                                  Expr::Classification ObjectClassification,
                                  ArrayRef<Expr *> Args,
                                  OverloadCandidateSet& CandidateSet,
                                  bool SuppressUserConversions,
                                  bool PartialOverloading) {
   if (!CandidateSet.isNewCandidate(MethodTmpl))
     return;
 
   // C++ [over.match.funcs]p7:
   //   In each case where a candidate is a function template, candidate
   //   function template specializations are generated using template argument
bool clang::isBetterOverloadCandidate(
   if (HasPS1 != HasPS2 && HasPS1)
     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
 /// declarations visible to be ambiguous in some cases (this happens when using
 /// a modularized libstdc++).
 bool Sema::isEquivalentInternalLinkageDeclaration(const NamedDecl *A,
                                                   const NamedDecl *B) {
   auto *VA = dyn_cast_or_null<ValueDecl>(A);
   auto *VB = dyn_cast_or_null<ValueDecl>(B);

lib/Sema/SemaType.cpp

case NullabilityKind::Nullable:
     return createSimpleAttr<TypeNullableAttr>(Ctx, Attr);
 
   case NullabilityKind::Unspecified:
     return createSimpleAttr<TypeNullUnspecifiedAttr>(Ctx, Attr);
   }
   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);
 
 static TypeSourceInfo *GetFullTypeForDeclarator(TypeProcessingState &state,
                                                 QualType declSpecType,
                                                 TypeSourceInfo *TInfo) {
   // The TypeSourceInfo that this function returns will not be a null type.
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) ||
+        auto IsClassMember = [&]() {
+          return (!state.getDeclarator().getCXXScopeSpec().isEmpty() &&
+                  state.getDeclarator()
+                          .getCXXScopeSpec()
+                          .getScopeRep()
+                          ->getKind() == NestedNameSpecifier::TypeSpec) ||
                  state.getDeclarator().getContext() ==
-                 DeclaratorContext::MemberContext)) {
-          LangAS CurAS = EPI.TypeQuals.getAddressSpace();
+                     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);
       }
       break;
     }
     case DeclaratorChunk::MemberPointer: {
       // The scope spec must refer to a class, or be dependent.
       CXXScopeSpec &SS = DeclType.Mem.Scope();
       QualType ClsType;
if (addrSpace > max) {
       Diag(AttrLoc, diag::err_attribute_address_space_too_high)
           << (unsigned)max.getZExtValue() << AddrSpace->getSourceRange();
       return QualType();
     }
 
     LangAS ASIdx =
         getLangASFromTargetAS(static_cast<unsigned>(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 Context.getAddrSpaceQualType(T, ASIdx);
   }
 
   // A check with similar intentions as checking if a type already has an
   // address space except for on a dependent types, basically if the
   // current type is already a DependentAddressSpaceType then its already
   // lined up to have another address space on it and we can't have
if (!T.isNull()) {
           Attr.getRange(), Ctx, Attr.getAttributeSpellingListIndex(),
           static_cast<unsigned>(T.getQualifiers().getAddressSpace()));
       Type = State.getAttributedType(ASAttr, T, T);
     } else {
       Attr.setInvalid();
     }
   } 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())) {
     }
 
     Type = S.Context.getAddrSpaceQualType(Type, ASIdx);
   }
 }
 
 /// Does this type have a "direct" ownership qualifier?  That is,
 /// is it written like "__strong id", as opposed to something like
if ( // Do not deduce addr space for function return type and function type,
       // Do not deduce addr space for types used to define a typedef and the
       // typedef itself, except the pointee type of a pointer type which is used
       // to define the typedef.
       (D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_typedef &&
        !IsPointee) ||
       // 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;
   // Put OpenCL automatic variable in private address space.
   // OpenCL v1.2 s6.5:
   // The default address space name for arguments to a function in a
   // program, or local variables of a function is __private. All function
   // arguments shall be in the __private address space.

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/address-space-of-this-class-scope.cl

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

test/SemaOpenCLCXX/address_space_overloading.cl

 // 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;
 };
 
 RetGlob foo(const __global int *);
 RetGen foo(const __generic int *);
 
 void kernel k() {
   __global int *ArgGlob;
   __generic int *ArgGen;
   __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}}
 }

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