[OpenCL] Add generic AS to 'this' pointer

[mikael]

Summary

Address spaces are cast into generic before invoking the constructor.

[rjmccall]

lib/CodeGen/CGCall.cpp

79

I would suggest taking this opportunity to set up the AST to support declaring methods in an arbitrary address space, so that you can just ask a CXXMethodDecl what address space it's in. You don't have to actually add language support for that — OpenCL C++ would simply change the it to the generic address space instead of the default — but I think that's the right technical approach for implementing this, as opposed to adding a bunch of OpenCL C++ -specific logic all over the compiler that just hardcodes a different address space.

4030

Always use the performAddrSpaceConversion target hook if there's a semantic address-space conversion required. But really, this doesn't seem like the right place to be doing this; it ought to happen higher up when we're adding the this argument to the call, either explicitly in IRGen or implicitly by expecting the object expression to already yield a value in the right address space.

I could definitely believe that we don't currently create CastExprs for simple qualification conversions of the object argument of a C++ method call, since (ignoring these address-space conversions) they're always trivial.

[Anastasia]

lib/CodeGen/CGCall.cpp

79

I quite like this idea. Apart from providing more clean implementation, it opens opportunities for solving several problems that I am trying to understand how to address. Specifically I am trying to find a way to 'overload' methods based on the address space of the object.

For example, if an object is created in the address space 1 then programmers should be able to provide a method to be used for objects in such address space for efficiency or even correctness issue.

The reasons I am looking at it is that currently C++ doesn't make much sense for address spaces, because we are removing them to generate just one implementation with generic/default address space. However,

• Not all address spaces can be converted to generic/default address space. Example in OpenCL is constant AS that can't be converted to any other.
• Higher performance can be achieved on some HW when using specific address spaces instead of default.

I was wondering if a method qualifier is a good language solution for this? For example in OpenCL we could write something like:

class foo
{
public:
  void bar() __private; // implies bar(__private foo*)
  void bar() __constant; // implies bar(__constant foo*)
};

I guess in C++ it can be done similarly:

class foo
{
public:
  void bar() __attribute__((address_space(1)));
  void bar() __attribute__((address_space(2)));
};

I would quite like to solve this generically, not just for OpenCL. I think a lot of implementation can be unified/reused then.

Without this address spaces seem pretty useless with C++ because they are just cast away to generic/default and no specific address space ends up at the AST level at all. This means implementation will have to rely on the optimizers to recover/deduce address spaces. But I would quite like to provide a way for the developers to manually tune the code for address spaces, just as it was done for OpenCL C.

Let me know if you have any thought/suggestions.

[rjmccall]

lib/CodeGen/CGCall.cpp

79

I was wondering if a method qualifier is a good language solution for this? For example in OpenCL we could write something like:

Yes, I think that's a very natural extension of C++'s method-qualification rules for const and volatile. Overloads would then be resolved based on which address space was the best match.

Now, to briefly take a holistic perspective on the language design, this feature would *strongly* benefit from a way to make a method templated over the address space of this. Unfortunately, I don't think that's reasonable to solve in a language extension; it's really something that needs core language work. That would be a pretty big leap in scope; that said, if you're interested in pursuing it, I'd be happy to share some thoughts on how it'd look, and I think there are several people in the Clang community who could help you with putting a proposal before the committee.

[mikael]

lib/CodeGen/CGCall.cpp

4030

Thanks for the input! It seems like the name performAddrSpaceConversion does not exist in the code-base though.

[rjmccall]

lib/CodeGen/CGCall.cpp

4030

Sorry, performAddrSpaceCast.

[Anastasia]

lib/CodeGen/CGCall.cpp

79

Thanks! Yes, I believe some template-based approach could make this a lot more usable. I would be really interested to learn more about your ideas. I would also quite like to see this properly in the spec. I will follow up with you in a separate thread.

[rjmccall]

lib/AST/DeclCXX.cpp

2186–2187

I feel like the right design now is for getTypeQualifiers() to return a Qualifiers and not just a mask.

That will also help identify other places that need to be updated to support address-space qualification.

lib/CodeGen/CGCall.cpp

71–72

Please update the comment to say "ignoring method CVR qualification" instead.

Also, maybe we should just eliminate this and use CXXMethodDecl::getThisType. If we want to preserve the micro-optimization of sharing CGFunctionInfos for method types that only differ by CV qualification, maybe we should add a parameter to getThisType to not add ABI-compatible qualification.

lib/CodeGen/CGClass.cpp

2020

This is adding an extra level of pointer, I think.

2023

...I'm not sure when getLangASFromTargetAS was added, but it shouldn't exist; you need to pass down the language address space of This, which in general will come from the pointer type. And then the AS != LangAS::Default check should just be comparing the constructor's address space to that address space, not specifically LangAS::Default.

Same comments apply to the code in commonEmitCXXMemberOrOperatorCall.

[mikael]

lib/Sema/SemaOverload.cpp

5254

Note that this change allowed me to revert some of the CG changes.

[rjmccall]

include/clang/AST/DeclCXX.h

2189

The address-space qualifier should get stored in the same way on the FunctionProtoType as the other method qualifiers (at least as far as the public interface of FPT goes). It's not the same thing as an ordinary type qualifier on the type, and trying to represent it that way will lead to a lot of problems downstream.

You'll probably also have to do some extra work in SemaType to make sure that the attribute gets applied in the right place.

You'll need to mess around with the representation of FPT to make this efficient, because we don't want to make every instance of FPT have to pay to store a full Qualifiers. I'd do it like this:

• Go ahead and add a full Qualifiers in FunctionProtoType::ExtProtoInfo; this is never stored permanently, so the cost is negligible.
• There are currently four bits for "qualifiers" in FunctionTypeBits. The only qualifiers that are actually important to store inline like this are const and volatile, but since we have four bits, let's store CVR here and use the fourth bit to indicate that there are more qualifiers stored in trailing storage.
• When there are non-CVR qualifiers present, ASTContext::getFunctionTypeInternal should allocate additional trailing storage for a Qualifiers object that can store the rest of the qualifiers.
• getTypeQuals() can check for the bit and load the qualifiers if present, or else just return the CVR qualifiers stored inline.

[Anastasia]

lib/Sema/SemaType.cpp

4817

Can this be moved into deduceOpenCLImplicitAddrSpace?

[mikael]

lib/Sema/SemaType.cpp

4817

deduceOpenCLImplicitAddrSpace is only wrapping a type with an address space. I also need to make sure the FunctionProtoType gets generated with the address space.

[Anastasia]

lib/Sema/SemaType.cpp

4817

FunctionProtoType is a type too. Does it go into deduceOpenCLImplicitAddrSpace when creating this?

[rjmccall]

I have a lot of comments about various things that need to be fixed, but I just want to take a moment to say that this is a tremendously impressive patch: given only some very high-level directions, you've done a great job figuring out what you needed to do to carry that out and make it work. Really well done.

lib/AST/ASTDumper.cpp

346 ↗

(On Diff #176596)

We must have some existing code to print an address space.

lib/AST/ItaniumMangle.cpp

1507

You can overload based on the address space, right? I think it needs to be mangled.

lib/AST/TypePrinter.cpp

804 ↗

(On Diff #176596)

You should be printing all the qualifiers here, including address spaces. If you want to suppress an implicit OpenCL address-space qualifier, that seems reasonable.

lib/CodeGen/CGDebugInfo.cpp

2593 ↗

(On Diff #176596)

I think this is just GetThisType again. Maybe there should be a version of that takes a record and an FPT; you could make it a static method on CXXMethodDecl. At any rate, I don't think you shouldn't be dropping the address space.

lib/Index/USRGeneration.cpp

274

Paging @akyrtzi here. The address-space qualifier should be part of the USR but I don't think if there's a defined schema for that.

lib/Parse/ParseCXXInlineMethods.cpp

419

This should take the full qualifiers. You can see how it ends up getting reassembled into the this type in the CXXThisScopeRAII ctor, and that type definitely needs to be address-space-qualified.

lib/Sema/SemaCodeComplete.cpp

1046

Hmm. This is actually going to break code completion for these methods, I think: IIUC in OpenCL every pointer and l-value is address-space-qualified, so only considering the CVRU qualifiers means that there will always be "extra" qualifiers on the pointer and code completion will think it doesn't match.

...I know code completion isn't your top priority, but this is probably easy enough to fix.

3701

This seems to be intentionally dropping non-CVR qualifiers, and I don't know why you ever would want to do that.

lib/Sema/SemaDeclCXX.cpp

11963

This should definitely not be dropping address spaces. You'll need to make a second patch to handle special members correctly, so you don't need to add a test for this yet, but you might as well get the logic right right now.

12332

Same thing as above.

lib/Sema/SemaOverload.cpp

1146

This is an algorithm that I think you need to get right and which shouldn't drop address spaces.

2828

This is fair to punt to a follow-up patch, but you *should* make that patch.

5070

You should be using the full qualifiers here.

lib/Sema/SemaTemplateDeduction.cpp

4660

Again, this is dropping non-CVR qualifiers for no particular reason.

[mikael]

Thanks for the feedback, I'll work on fixing the issues!

lib/Sema/SemaOverload.cpp

1146

But with this patch we can only have __generic address space for methods. So if I try to implement something here I can't really test it? Or am I missing something?

[Anastasia]

lib/AST/ItaniumMangle.cpp

1507

Does this refer to our earlier discussion https://reviews.llvm.org/D54862#inline-484509

We don't have a way to qualify methods with an address space yet? I was going to send an RFC to cfe-dev for this but if you think it would be ok to go ahead with an implementation, I am happy with it. Either way would it be better to do this in a separate patch?

[rjmccall]

lib/AST/ItaniumMangle.cpp

1507

I'm fine with delaying implementation on these two issues until a later patch since, as you say, they can't be tested well until we support arbitrary address-space qualifiers. Please at least leave FIXMEs for them.

[Anastasia]

lib/Sema/SemaOverload.cpp

1146

Since we didn't add address spaces to the method qualifiers yet, it seems logical to ignore them.

Although we could already say, since we would like to add such overloading anyway, it would be logical to return true for the address space qualifiers mismatch just like it's done for CV quals in line 1154. However, I can't think of any way we could test this indeed.

[mikael]

I uploaded a new patch for most of the comments. I did leave some things out since they need clarification.

[mikael]

lib/AST/ItaniumMangle.cpp

1507

I actually did update this, since we can test the mangling of __generic in the test.

[akyrtzi]

lib/Index/USRGeneration.cpp

274

If I understand correctly from other comments you can't add special mangling and USR generation yet and intend to add FIXME for doing mangling support later ? If yes, could you also add FIXME here and re-visit ?

[rjmccall]

lib/AST/ItaniumMangle.cpp

1507

Great, thanks. Looks good.

lib/AST/Type.cpp

2950 ↗

(On Diff #176831)

The indentation here is messed up.

You seem to be mixing up "fast qualifiers" with "CVR qualifiers" in a few places. That's currently correct, in the sense that the fast qualifiers are currently defined to be the CVR qualifiers, but the point of the distinction is that we might want to change that (and we have in fact explored that in the past). I think you should make FunctionTypeBits only claim to store the fast qualifiers, which mostly just means updating a few field / accessor names and changing things like the getCVRQualifiers() call right above this to be getFastQualifiers().

If there are places where it's convenient to assume that the CVR qualifiers are exactly the fast qualifiers, it's okay to static_assert that, but you should make sure to static_assert it in each place you assume it.

lib/CodeGen/CGDebugInfo.cpp

2593 ↗

(On Diff #176831)

Why is this getMostRecentCXXRecordDecl instead of getClass?

lib/CodeGen/CGDeclCXX.cpp

32

What causes OpenCL C++ declarations to go down this path?

lib/Index/USRGeneration.cpp

274

While you're here, can you described the right way for us to extend USR generation here to potentially add an address space?

lib/Parse/ParseDecl.cpp

6166 ↗

(On Diff #176831)

I think fromCVRUQualifiers is probably the right function, and if that doesn't exist you should add it.

lib/Sema/SemaExprCXX.cpp

1112 ↗

(On Diff #176831)

I don't think there's a good reason for this to be dropping non-CVR qualifiers. If there is, it really needs a comment.

lib/Sema/TreeTransform.h

5276 ↗

(On Diff #176831)

I think you need to turn ThisTypeQuals into a Qualifiers here.

[mikael]

lib/AST/Type.cpp

2950 ↗

(On Diff #176831)

I'll change it, but I do have a question related to this:

Can we make any assumptions with the "fast qualifiers"? I'd like it if we can assume that they fit in 3 bits. Otherwise I think I also need an assertion here.

lib/CodeGen/CGDebugInfo.cpp

2593 ↗

(On Diff #176831)

Your initial comment suggested to send in a record rather than a type. But I see now that it may make more sense to use the type directly instead. Will update it.

lib/CodeGen/CGDeclCXX.cpp

32

It goes there when we defined an object in __local address space. From the spec "The constructors of a objects in local memory should only be initialized once". This is the same case as with C++ 11 static local variables.

[mikael]

• Added FIXME's.
• Fixed the latest comments.

[mikael]

lib/AST/Type.cpp

2950 ↗

(On Diff #176831)

I solved it like this in the end:

• I defined the FastTypeQuals to get the size of Qualifiers::FastWidth to make it dependent on possible changes to the fast flags.
• Then I put a static_assertion to guard the hastConst(), hasVolatile() and hasRestrict() methods.

[rjmccall]

LGTM, but you can probably undo one of my requests. :)

lib/AST/Type.cpp

2950 ↗

(On Diff #176831)

Great, looks good.

lib/CodeGen/CGDebugInfo.cpp

2593 ↗

(On Diff #176831)

Oh, I actually didn't even think about the fact that the getClass()) returns a Type* instead of a CXXRecordDecl. Using the record decl is probably better, sorry.

[akyrtzi]

lib/Index/USRGeneration.cpp

274

I'm not familiar with the mechanism, is address space supposed to be identified via a single integer ? If so, then I think adding '-' character plus the integer as character for the address space, immediately after the type qualifiers, should be a way to go.

[mikael]

I rebased on Friday, and noticed that I broke two tests:

Failing Tests (2):

Clang :: CodeGenOpenCLCXX/template-address-spaces.cl
Clang :: SemaOpenCLCXX/address-space-templates.cl

This upload contains a few extra fixes.

[mikael]

lib/Sema/SemaDecl.cpp

3196

When merging the class function and the file context function the address space was lost here.

lib/Sema/SemaType.cpp

4818

The first issue was that if a function is declared outside of a class this code did not deduce it. Now it does.

lib/Sema/TreeTransform.h

4277 ↗

(On Diff #177477)

Another place where the address space was removed.

test/SemaOpenCLCXX/address-space-templates.cl

7 ↗

(On Diff #177477)

This was the remaining issue that I have not solved yet. It looked like this issue was not so trivial so I think it makes sense to postpone it.

[Anastasia]

lib/Sema/SemaType.cpp

3922

From the state you can use getDeclarator() that then call a method getContext().

5127

Rather than passing DeclContext you can get it from the Declarator using getContext() method.

lib/Sema/TreeTransform.h

4277 ↗

(On Diff #177477)

Why OpenCL in the comment, since the code seems to be rather generic.

Also not sure this comments adds anything that can't be understood from the code you are adding.

test/SemaOpenCLCXX/address-space-templates.cl

7 ↗

(On Diff #177477)

Do you know why it appears twice?

[mikael]

test/SemaOpenCLCXX/address-space-templates.cl

7 ↗

(On Diff #177477)

I believe the main issue is that we have some code that does the following in SemaTemplateInstantiateDecl.cpp:

OldTL.getAs<FunctionProtoTypeLoc>()

But the implementation of getAs returns nullptr if you have local modifiers. So since this code will always return nullptr we end up taking a different paths.

[mikael]

lib/Sema/SemaType.cpp

5127

Declarator::getContext() returns an enumeration DeclaratorContext which is not enough in this case.

[Anastasia]

LGTM! Apart from small improvement that can done before committing if it works at all. :)

lib/Sema/SemaType.cpp

5127

Ok, I was wondering if you can use scope clang::NestedNameSpecifier::SpecifierKind to check if something is a method instead of passing DC down.

[mikael]

Seems like my this commit broke: http://lab.llvm.org:8011/builders/lldb-x86_64-ubuntu-14.04-buildserver/builds/33176

But since I don't really know what anything about lldb, I probably won't be able to fix it fast enough.

/lldb-buildbot/buildServerSlave/lldb-android-buildserver/llvm/tools/lldb/source/Symbol/ClangASTContext.cpp:2208:24: error: no viable overloaded '='
  proto_info.TypeQuals = type_quals;
  ~~~~~~~~~~~~~~~~~~~~ ^ ~~~~~~~~~~
/lldb-buildbot/buildServerSlave/lldb-android-buildserver/llvm/tools/clang/include/clang/AST/Type.h:141:7: note: candidate function (the implicit copy assignment operator) not viable: no known conversion from 'unsigned int' to 'const clang::Qualifiers' for 1st argument
class Qualifiers {
      ^
/lldb-buildbot/buildServerSlave/lldb-android-buildserver/llvm/tools/clang/include/clang/AST/Type.h:141:7: note: candidate function (the implicit move assignment operator) not viable: no known conversion from 'unsigned int' to 'clang::Qualifiers' for 1st argument
1 error generated.

I think it can be easily solved by Qualifiers::fromFastMask(type_quals); also updating the variable name since it changed.

But I know nothing about lldb, so I better revert this patch for now.

[riccibruno]

include/clang/AST/Type.h

3692

You can use FunctionType::FunctionTypeExtraBitfields
to store the qualifiers. It was designed to store extra data.
No need to add yet another trailing object.

[rjmccall]

In D54862#1328290, @mikael wrote:

Seems like my this commit broke: http://lab.llvm.org:8011/builders/lldb-x86_64-ubuntu-14.04-buildserver/builds/33176

But since I don't really know what anything about lldb, I probably won't be able to fix it fast enough.

/lldb-buildbot/buildServerSlave/lldb-android-buildserver/llvm/tools/lldb/source/Symbol/ClangASTContext.cpp:2208:24: error: no viable overloaded '='
  proto_info.TypeQuals = type_quals;
  ~~~~~~~~~~~~~~~~~~~~ ^ ~~~~~~~~~~
/lldb-buildbot/buildServerSlave/lldb-android-buildserver/llvm/tools/clang/include/clang/AST/Type.h:141:7: note: candidate function (the implicit copy assignment operator) not viable: no known conversion from 'unsigned int' to 'const clang::Qualifiers' for 1st argument
class Qualifiers {
      ^
/lldb-buildbot/buildServerSlave/lldb-android-buildserver/llvm/tools/clang/include/clang/AST/Type.h:141:7: note: candidate function (the implicit move assignment operator) not viable: no known conversion from 'unsigned int' to 'clang::Qualifiers' for 1st argument
1 error generated.

I think it can be easily solved by Qualifiers::fromFastMask(type_quals); also updating the variable name since it changed.

I think that's exactly what you need to do, and you'd be welcome to commit that instead of reverting this.

lib/Index/USRGeneration.cpp

274

Okay. Some of them are integers; others are identified in the language as more of a keyword and so there's no stable mapping from an integer value. Should we just come up with some simple schema for that following - that's unambiguous with a digit?

[ebevhan]

lib/CodeGen/CGCall.cpp

75

I'm a bit late to the party, but shouldn't this be MD->getTypeQualifiers().getAddressSpace()?

[rjmccall]

lib/CodeGen/CGCall.cpp

75

Yes, you're right, it should.

[ebevhan]

I'm also a bit confused about the semantics that this patch is applying to function types. It mostly seems to concern the extra trailing Qualifiers on CXXMethodDecl to store the addrspace quals, but in some places (SemaType:4842, SemaDecl:3198) it seems to be applying the address space to the function type itself, which certainly seems like something else to me. A function with an address space qualified type would (to me, at least) be a function *located* in that address space, not one qualified to take a this from that address space.

[rjmccall]

In D54862#1333700, @ebevhan wrote:

I'm also a bit confused about the semantics that this patch is applying to function types. It mostly seems to concern the extra trailing Qualifiers on CXXMethodDecl to store the addrspace quals, but in some places (SemaType:4842, SemaDecl:3198) it seems to be applying the address space to the function type itself, which certainly seems like something else to me. A function with an address space qualified type would (to me, at least) be a function *located* in that address space, not one qualified to take a this from that address space.

Yeah, there may be some confusion about that, and unfortunately it's easy to miss in the review.

But yes, I agree that address-space qualifiers are in principle one of the few qualifiers that make sense to be able to have (in a normal way) on function types.

include/clang/AST/DeclCXX.h

public:
   /// Returns the type of the \c this pointer.
   ///
   /// Should only be called for instance (i.e., non-static) methods. Note
   /// that for the call operator of a lambda closure type, this returns the
   /// desugared 'this' type (a pointer to the closure type), not the captured
   /// 'this' type.
   QualType getThisType(ASTContext &C) const;
 
-  unsigned getTypeQualifiers() const {
-    return getType()->getAs<FunctionProtoType>()->getTypeQuals();
+  Qualifiers getTypeQualifiers() const {
+    Qualifiers C;
+    unsigned CVRU = getType()->getAs<FunctionProtoType>()->getTypeQuals();
+    C.addCVRUQualifiers(CVRU);
+    C.addQualifiers(getType().getQualifiers());
+    return C;
   }
 
   /// Retrieve the ref-qualifier associated with this method.
   ///
   /// In the following example, \c f() has an lvalue ref-qualifier, \c g()
   /// has an rvalue ref-qualifier, and \c h() has no ref-qualifier.
   /// @code
   /// struct X {

include/clang/AST/Type.h

public:
   void setRestrict(bool flag) {
     Mask = (Mask & ~Restrict) | (flag ? Restrict : 0);
   }
   void removeRestrict() { Mask &= ~Restrict; }
   void addRestrict() { Mask |= Restrict; }
 
   bool hasCVRQualifiers() const { return getCVRQualifiers(); }
   unsigned getCVRQualifiers() const { return Mask & CVRMask; }
+  unsigned getCVRUQualifiers() const { return Mask & (CVRMask | UMask); }
+
   void setCVRQualifiers(unsigned mask) {
     assert(!(mask & ~CVRMask) && "bitmask contains non-CVR bits");
     Mask = (Mask & ~CVRMask) | mask;
   }
   void removeCVRQualifiers(unsigned mask) {
     assert(!(mask & ~CVRMask) && "bitmask contains non-CVR bits");
     Mask &= ~mask;
   }

 /// which optional. For more information about the trailing objects see
 /// the first comment inside FunctionProtoType.
 class FunctionProtoType final
     : public FunctionType,
       public llvm::FoldingSetNode,
       private llvm::TrailingObjects<
           FunctionProtoType, QualType, FunctionType::FunctionTypeExtraBitfields,
           FunctionType::ExceptionType, Expr *, FunctionDecl *,
           FunctionType::ExtParameterInfo> {
   friend class ASTContext; // ASTContext creates these.
   friend TrailingObjects;
 
   // FunctionProtoType is followed by several trailing objects, some of
   // which optional. They are in order:
   //
   // * An array of getNumParams() QualType holding the parameter types.
   //   Always present. Note that for the vast majority of FunctionProtoType,

lib/AST/DeclCXX.cpp

QualType CXXMethodDecl::getThisType(ASTContext &C) const {
   // If the member function is declared const, the type of this is const X*,
   // if the member function is declared volatile, the type of this is
   // volatile X*, and if the member function is declared const volatile,
   // the type of this is const volatile X*.
 
   assert(isInstance() && "No 'this' for static methods!");
 
   QualType ClassTy = C.getTypeDeclType(getParent());
-  ClassTy = C.getQualifiedType(ClassTy,
-                               Qualifiers::fromCVRUMask(getTypeQualifiers()));
+  ClassTy = C.getQualifiedType(ClassTy, getTypeQualifiers());
+
   return C.getPointerType(ClassTy);
 }
 
 bool CXXMethodDecl::hasInlineBody() const {
   // If this function is a template instantiation, look at the template from
   // which it was instantiated.
   const FunctionDecl *CheckFn = getTemplateInstantiationPattern();
   if (!CheckFn)

lib/AST/ItaniumMangle.cpp

void CXXNameMangler::mangleNestedName(const NamedDecl *ND,
                                       bool NoFunction) {
   // <nested-name>
   //   ::= N [<CV-qualifiers>] [<ref-qualifier>] <prefix> <unqualified-name> E
   //   ::= N [<CV-qualifiers>] [<ref-qualifier>] <template-prefix>
   //       <template-args> E
 
   Out << 'N';
   if (const CXXMethodDecl *Method = dyn_cast<CXXMethodDecl>(ND)) {
-    Qualifiers MethodQuals =
-        Qualifiers::fromCVRUMask(Method->getTypeQualifiers());
+    Qualifiers MethodQuals = Qualifiers::fromCVRUMask(
+        Method->getTypeQualifiers().getCVRUQualifiers());
     // We do not consider restrict a distinguishing attribute for overloading
     // purposes so we must not mangle it.
     MethodQuals.removeRestrict();
     mangleQualifiers(MethodQuals);
     mangleRefQualifier(Method->getRefQualifier());
   }
 
   // Check if we have a template.

lib/CodeGen/CGCall.cpp

unsigned CodeGenTypes::ClangCallConvToLLVMCallConv(CallingConv CC) {
   case CC_SpirFunction: return llvm::CallingConv::SPIR_FUNC;
   case CC_OpenCLKernel: return CGM.getTargetCodeGenInfo().getOpenCLKernelCallingConv();
   case CC_PreserveMost: return llvm::CallingConv::PreserveMost;
   case CC_PreserveAll: return llvm::CallingConv::PreserveAll;
   case CC_Swift: return llvm::CallingConv::Swift;
   }
 }
 
-/// Derives the 'this' type for codegen purposes, i.e. ignoring method
+/// Derives the 'this' type for codegen purposes, i.e. ignoring method CVR
 /// qualification.
-/// FIXME: address space qualification?
-static CanQualType GetThisType(ASTContext &Context, const CXXRecordDecl *RD) {
+static CanQualType GetThisType(ASTContext &Context, const CXXRecordDecl *RD, const CXXMethodDecl *MD) {
   QualType RecTy = Context.getTagDeclType(RD)->getCanonicalTypeInternal();
+  if (MD)
+    RecTy = Context.getAddrSpaceQualType(RecTy, MD->getType().getAddressSpace());
   return Context.getPointerType(CanQualType::CreateUnsafe(RecTy));
 }
 
 /// Returns the canonical formal type of the given C++ method.
 static CanQual<FunctionProtoType> GetFormalType(const CXXMethodDecl *MD) {
   return MD->getType()->getCanonicalTypeUnqualified()
            .getAs<FunctionProtoType>();
 }
 
 /// Returns the "extra-canonicalized" return type, which discards
 /// qualifiers on the return type.  Codegen doesn't care about them,
 /// and it makes ABI code a little easier to be able to assume that

 const CGFunctionInfo &
 CodeGenTypes::arrangeCXXMethodType(const CXXRecordDecl *RD,
                                    const FunctionProtoType *FTP,
                                    const CXXMethodDecl *MD) {
   SmallVector<CanQualType, 16> argTypes;
 
   // Add the 'this' pointer.
   if (RD)
-    argTypes.push_back(GetThisType(Context, RD));
+    argTypes.push_back(GetThisType(Context, RD, MD));
   else
     argTypes.push_back(Context.VoidPtrTy);
 
   return ::arrangeLLVMFunctionInfo(
       *this, true, argTypes,
       FTP->getCanonicalTypeUnqualified().getAs<FunctionProtoType>(), MD);
 }
 
bool CodeGenTypes::inheritingCtorHasParams(
   }
 
 const CGFunctionInfo &
 CodeGenTypes::arrangeCXXStructorDeclaration(const CXXMethodDecl *MD,
                                             StructorType Type) {
 
   SmallVector<CanQualType, 16> argTypes;
   SmallVector<FunctionProtoType::ExtParameterInfo, 16> paramInfos;
-  argTypes.push_back(GetThisType(Context, MD->getParent()));
+  argTypes.push_back(GetThisType(Context, MD->getParent(), MD));
 
   bool PassParams = true;
 
   GlobalDecl GD;
   if (auto *CD = dyn_cast<CXXConstructorDecl>(MD)) {
     GD = GlobalDecl(CD, toCXXCtorType(Type));
 
     // A base class inheriting constructor doesn't get forwarded arguments

 /// varargs.  Return a void pointer, regardless of the actual return type.
 /// The body of the thunk will end in a musttail call to a function of the
 /// correct type, and the caller will bitcast the function to the correct
 /// prototype.
 const CGFunctionInfo &
 CodeGenTypes::arrangeUnprototypedMustTailThunk(const CXXMethodDecl *MD) {
   assert(MD->isVirtual() && "only methods have thunks");
   CanQual<FunctionProtoType> FTP = GetFormalType(MD);
-  CanQualType ArgTys[] = { GetThisType(Context, MD->getParent()) };
+  CanQualType ArgTys[] = { GetThisType(Context, MD->getParent(), MD) };
   return arrangeLLVMFunctionInfo(Context.VoidTy, /*instanceMethod=*/false,
                                  /*chainCall=*/false, ArgTys,
                                  FTP->getExtInfo(), {}, RequiredArgs(1));
 }
 
 const CGFunctionInfo &
 CodeGenTypes::arrangeMSCtorClosure(const CXXConstructorDecl *CD,
                                    CXXCtorType CT) {
   assert(CT == Ctor_CopyingClosure || CT == Ctor_DefaultClosure);
 
   CanQual<FunctionProtoType> FTP = GetFormalType(CD);
   SmallVector<CanQualType, 2> ArgTys;
   const CXXRecordDecl *RD = CD->getParent();
-  ArgTys.push_back(GetThisType(Context, RD));
+  ArgTys.push_back(GetThisType(Context, RD, CD));
   if (CT == Ctor_CopyingClosure)
     ArgTys.push_back(*FTP->param_type_begin());
   if (RD->getNumVBases() > 0)
     ArgTys.push_back(Context.IntTy);
   CallingConv CC = Context.getDefaultCallingConvention(
       /*IsVariadic=*/false, /*IsCXXMethod=*/true);
   return arrangeLLVMFunctionInfo(Context.VoidTy, /*instanceMethod=*/true,
                                  /*chainCall=*/false, ArgTys,
case ABIArgInfo::Direct: {
 
         // If the argument doesn't match, perform a bitcast to coerce it.  This
         // can happen due to trivial type mismatches.
         if (FirstIRArg < IRFuncTy->getNumParams() &&
             V->getType() != IRFuncTy->getParamType(FirstIRArg))
           V = Builder.CreateBitCast(V, IRFuncTy->getParamType(FirstIRArg));
 
         IRCallArgs[FirstIRArg] = V;
         break;
       }
 
       // FIXME: Avoid the conversion through memory if possible.
       Address Src = Address::invalid();
       if (!I->isAggregate()) {
         Src = CreateMemTemp(I->Ty, "coerce");
         I->copyInto(*this, Src);
       } else {

lib/CodeGen/CGClass.cpp

 //
 //===----------------------------------------------------------------------===//
 
 #include "CGBlocks.h"
 #include "CGCXXABI.h"
 #include "CGDebugInfo.h"
 #include "CGRecordLayout.h"
 #include "CodeGenFunction.h"
+#include "TargetInfo.h"
 #include "clang/AST/CXXInheritance.h"
 #include "clang/AST/DeclTemplate.h"
 #include "clang/AST/EvaluatedExprVisitor.h"
 #include "clang/AST/RecordLayout.h"
 #include "clang/AST/StmtCXX.h"
 #include "clang/Basic/TargetBuiltins.h"
 #include "clang/CodeGen/CGFunctionInfo.h"
 #include "clang/Frontend/CodeGenOptions.h"
void CodeGenFunction::EmitCXXConstructorCall(const CXXConstructorDecl *D,
                                              CXXCtorType Type,
                                              bool ForVirtualBase,
                                              bool Delegating, Address This,
                                              const CXXConstructExpr *E,
                                              AggValueSlot::Overlap_t Overlap,
                                              bool NewPointerIsChecked) {
   CallArgList Args;
 
+  LangAS SlotAS = E->getType().getAddressSpace();
+  QualType ThisType = D->getThisType(getContext());
+  LangAS ThisAS = ThisType.getTypePtr()->getPointeeType().getAddressSpace();
+  llvm::Value *ThisPtr = This.getPointer();
+  if (SlotAS != ThisAS) {
+    unsigned TargetThisAS = getContext().getTargetAddressSpace(ThisAS);
+    llvm::Type *NewType =
+        ThisPtr->getType()->getPointerElementType()->getPointerTo(TargetThisAS);
+    ThisPtr = getTargetHooks().performAddrSpaceCast(*this, This.getPointer(),
+                                                    ThisAS, SlotAS, NewType);
+  }
   // Push the this ptr.
-  Args.add(RValue::get(This.getPointer()), D->getThisType(getContext()));
+  Args.add(RValue::get(ThisPtr), D->getThisType(getContext()));
 
   // If this is a trivial constructor, emit a memcpy now before we lose
   // the alignment information on the argument.
   // FIXME: It would be better to preserve alignment information into CallArg.
   if (isMemcpyEquivalentSpecialMember(D)) {
     assert(E->getNumArgs() == 1 && "unexpected argcount for trivial ctor");
 
     const Expr *Arg = E->getArg(0);

lib/CodeGen/CGDeclCXX.cpp

 #include "llvm/IR/MDBuilder.h"
 #include "llvm/Support/Path.h"
 
 using namespace clang;
 using namespace CodeGen;
 
 static void EmitDeclInit(CodeGenFunction &CGF, const VarDecl &D,
                          ConstantAddress DeclPtr) {
-  assert(D.hasGlobalStorage() && "VarDecl must have global storage!");
+  assert(
+      (D.hasGlobalStorage() ||
+       (D.hasLocalStorage() && CGF.getContext().getLangOpts().OpenCLCPlusPlus)) &&
+      "VarDecl must have global or local (in the case of OpenCL) storage!");
   assert(!D.getType()->isReferenceType() &&
          "Should not call EmitDeclInit on a reference!");
 
   QualType type = D.getType();
   LValue lv = CGF.MakeAddrLValue(DeclPtr, type);
 
   const Expr *Init = D.getInit();
   switch (CGF.getEvaluationKind(type)) {

lib/Index/USRGeneration.cpp

if (IsTemplate) {
     // \endcode
     Out << '#';
     VisitType(D->getReturnType());
   }
   Out << '#';
   if (const CXXMethodDecl *MD = dyn_cast<CXXMethodDecl>(D)) {
     if (MD->isStatic())
       Out << 'S';
-    if (unsigned quals = MD->getTypeQualifiers())
+    if (unsigned quals = MD->getTypeQualifiers().getCVRUQualifiers())
       Out << (char)('0' + quals);
     switch (MD->getRefQualifier()) {
     case RQ_None: break;
     case RQ_LValue: Out << '&'; break;
     case RQ_RValue: Out << "&&"; break;
     }
   }
 }
 

lib/Parse/ParseCXXInlineMethods.cpp

if (CachedTokens *Toks = LM.ExceptionSpecTokens) {
     CXXMethodDecl *Method;
     if (FunctionTemplateDecl *FunTmpl
           = dyn_cast<FunctionTemplateDecl>(LM.Method))
       Method = cast<CXXMethodDecl>(FunTmpl->getTemplatedDecl());
     else
       Method = cast<CXXMethodDecl>(LM.Method);
 
     Sema::CXXThisScopeRAII ThisScope(Actions, Method->getParent(),
-                                     Method->getTypeQualifiers(),
+                                     Method->getTypeQualifiers().getCVRUQualifiers(),
                                      getLangOpts().CPlusPlus11);
 
     // Parse the exception-specification.
     SourceRange SpecificationRange;
     SmallVector<ParsedType, 4> DynamicExceptions;
     SmallVector<SourceRange, 4> DynamicExceptionRanges;
     ExprResult NoexceptExpr;
     CachedTokens *ExceptionSpecTokens;

lib/Sema/SemaCodeComplete.cpp

void ResultBuilder::AddResult(Result R, DeclContext *CurContext,
   if (InBaseClass)
     setInBaseClass(R);
 
   AdjustResultPriorityForDecl(R);
 
   if (HasObjectTypeQualifiers)
     if (const CXXMethodDecl *Method = dyn_cast<CXXMethodDecl>(R.Declaration))
       if (Method->isInstance()) {
-        Qualifiers MethodQuals
-                        = Qualifiers::fromCVRMask(Method->getTypeQualifiers());
+        Qualifiers MethodQuals = Qualifiers::fromCVRMask(
+            Method->getTypeQualifiers().getCVRQualifiers());
         if (ObjectTypeQualifiers == MethodQuals)
           R.Priority += CCD_ObjectQualifierMatch;
         else if (ObjectTypeQualifiers - MethodQuals) {
           // The method cannot be invoked, because doing so would drop
           // qualifiers.
           return;
         }
       }
void Sema::CodeCompleteOrdinaryName(Scope *S,
 
   // If we are in a C++ non-static member function, check the qualifiers on
   // the member function to filter/prioritize the results list and set the
   // context to the record context so that accessibility check in base class
   // works correctly.
   RecordDecl *MemberCompletionRecord = nullptr;
   if (CXXMethodDecl *CurMethod = dyn_cast<CXXMethodDecl>(CurContext)) {
     if (CurMethod->isInstance()) {
-      Results.setObjectTypeQualifiers(
-                      Qualifiers::fromCVRMask(CurMethod->getTypeQualifiers()));
+      Results.setObjectTypeQualifiers(Qualifiers::fromCVRMask(
+          CurMethod->getTypeQualifiers().getCVRQualifiers()));
       MemberCompletionRecord = CurMethod->getParent();
     }
   }
 
   CodeCompletionDeclConsumer Consumer(Results, CurContext, /*FixIts=*/{},
                                       MemberCompletionRecord);
   LookupVisibleDecls(S, LookupOrdinaryName, Consumer,
                      CodeCompleter->includeGlobals(),

lib/Sema/SemaDecl.cpp

if (OldTypeInfo.getNoCallerSavedRegs() !=
     NewTypeInfo = NewTypeInfo.withNoCallerSavedRegs(true);
     RequiresAdjustment = true;
   }
 
   if (RequiresAdjustment) {
     const FunctionType *AdjustedType = New->getType()->getAs<FunctionType>();
     AdjustedType = Context.adjustFunctionType(AdjustedType, NewTypeInfo);
     New->setType(QualType(AdjustedType, 0));
     NewQType = Context.getCanonicalType(New->getType());
     NewType = cast<FunctionType>(NewQType);
   }
 
   // If this redeclaration makes the function inline, we may need to add it to
   // UndefinedButUsed.
   if (!Old->isInlined() && New->isInlined() &&
       !New->hasAttr<GNUInlineAttr>() &&
       !getLangOpts().GNUInline &&
bool Sema::CheckFunctionDeclaration(Scope *S, FunctionDecl *NewFD,
   // This needs to be delayed until we know whether this is an out-of-line
   // definition of a static member function.
   //
   // This rule is not present in C++1y, so we produce a backwards
   // compatibility warning whenever it happens in C++11.
   CXXMethodDecl *MD = dyn_cast<CXXMethodDecl>(NewFD);
   if (!getLangOpts().CPlusPlus14 && MD && MD->isConstexpr() &&
       !MD->isStatic() && !isa<CXXConstructorDecl>(MD) &&
-      (MD->getTypeQualifiers() & Qualifiers::Const) == 0) {
+      !MD->getTypeQualifiers().hasConst()) {
     CXXMethodDecl *OldMD = nullptr;
     if (OldDecl)
       OldMD = dyn_cast_or_null<CXXMethodDecl>(OldDecl->getAsFunction());
     if (!OldMD || !OldMD->isStatic()) {
       const FunctionProtoType *FPT =
         MD->getType()->castAs<FunctionProtoType>();
       FunctionProtoType::ExtProtoInfo EPI = FPT->getExtProtoInfo();
       EPI.TypeQuals |= Qualifiers::Const;

lib/Sema/SemaDeclCXX.cpp

for (auto &Base : ClassDecl->bases()) {
 
     // Construct the "from" expression, which is an implicit cast to the
     // appropriately-qualified base type.
     CastBuilder From(OtherRef, Context.getQualifiedType(BaseType, OtherQuals),
                      VK_LValue, BasePath);
 
     // Dereference "this".
     DerefBuilder DerefThis(This);
-    CastBuilder To(DerefThis,
-                   Context.getCVRQualifiedType(
-                       BaseType, CopyAssignOperator->getTypeQualifiers()),
-                   VK_LValue, BasePath);
+    CastBuilder To(
+        DerefThis,
+            BaseType,
+            CopyAssignOperator->getTypeQualifiers().getCVRUQualifiers()),
 
     // Build the copy.
     StmtResult Copy = buildSingleCopyAssign(*this, Loc, BaseType,
                                             To, From,
                                             /*CopyingBaseSubobject=*/true,
                                             /*Copying=*/true);
     if (Copy.isInvalid()) {
       CopyAssignOperator->setInvalidDecl();
for (auto &Base : ClassDecl->bases()) {
     // Construct the "from" expression, which is an implicit cast to the
     // appropriately-qualified base type.
     CastBuilder From(OtherRef, BaseType, VK_XValue, BasePath);
 
     // Dereference "this".
     DerefBuilder DerefThis(This);
 
     // Implicitly cast "this" to the appropriately-qualified base type.
-    CastBuilder To(DerefThis,
-                   Context.getCVRQualifiedType(
-                       BaseType, MoveAssignOperator->getTypeQualifiers()),
-                   VK_LValue, BasePath);
+    CastBuilder To(
+        DerefThis,
+            BaseType,
+            MoveAssignOperator->getTypeQualifiers().getCVRQualifiers()),
 
     // Build the move.
     StmtResult Move = buildSingleCopyAssign(*this, Loc, BaseType,
                                             To, From,
                                             /*CopyingBaseSubobject=*/true,
                                             /*Copying=*/false);
     if (Move.isInvalid()) {
       MoveAssignOperator->setInvalidDecl();

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.
-    unsigned OldQuals = OldMethod->getTypeQualifiers();
-    unsigned NewQuals = NewMethod->getTypeQualifiers();
+    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)
if (!Context.hasSameType(FromFunction->getReturnType(),
                            ToFunction->getReturnType())) {
     PDiag << ft_return_type << ToFunction->getReturnType()
           << FromFunction->getReturnType();
     return;
   }
 
   unsigned FromQuals = FromFunction->getTypeQuals(),
            ToQuals = ToFunction->getTypeQuals();
   if (FromQuals != ToQuals) {
     PDiag << ft_qualifer_mismatch << ToQuals << FromQuals;
     return;
   }
 
   // Handle exception specification differences on canonical type (in C++17
   // onwards).
   if (cast<FunctionProtoType>(FromFunction->getCanonicalTypeUnqualified())
           ->isNothrow() !=

 static ImplicitConversionSequence
 TryObjectArgumentInitialization(Sema &S, SourceLocation Loc, QualType FromType,
                                 Expr::Classification FromClassification,
                                 CXXMethodDecl *Method,
                                 CXXRecordDecl *ActingContext) {
   QualType ClassType = S.Context.getTypeDeclType(ActingContext);
   // [class.dtor]p2: A destructor can be invoked for a const, volatile or
   //                 const volatile object.
-  unsigned Quals = isa<CXXDestructorDecl>(Method) ?
-    Qualifiers::Const | Qualifiers::Volatile : Method->getTypeQualifiers();
+  unsigned Quals = isa<CXXDestructorDecl>(Method)
+                       ? Qualifiers::Const | Qualifiers::Volatile
+                       : Method->getTypeQualifiers().getCVRUQualifiers();
   QualType ImplicitParamType =  S.Context.getCVRQualifiedType(ClassType, Quals);
 
   // Set up the conversion sequence as a "bad" conversion, to allow us
   // to exit early.
   ImplicitConversionSequence ICS;
 
   // We need to have an object of class type.
   if (const PointerType *PT = FromType->getAs<PointerType>()) {
Sema::PerformObjectArgumentInitialization(Expr *From,
   if (ICS.Standard.Second == ICK_Derived_To_Base) {
     ExprResult FromRes =
       PerformObjectMemberConversion(From, Qualifier, FoundDecl, Method);
     if (FromRes.isInvalid())
       return ExprError();
     From = FromRes.get();
   }
 
-  if (!Context.hasSameType(From->getType(), DestType))
+  if (!Context.hasSameType(From->getType(), DestType)) {
+    if (From->getType().getAddressSpace() != DestType.getAddressSpace())
+      From = ImpCastExprToType(From, DestType, CK_AddressSpaceConversion,
+                             From->getValueKind()).get();
+    else
       From = ImpCastExprToType(From, DestType, CK_NoOp,
                              From->getValueKind()).get();
+  }
   return From;
 }
 
 /// TryContextuallyConvertToBool - Attempt to contextually convert the
 /// expression From to bool (C++0x [conv]p3).
 static ImplicitConversionSequence
 TryContextuallyConvertToBool(Sema &S, Expr *From) {
   return TryImplicitConversion(S, From, S.Context.BoolTy,

lib/Sema/SemaTemplateDeduction.cpp

QualType ResultType;
     //   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
     //   and the end of the function-definition, member-declarator, or
     //   declarator.
     unsigned ThisTypeQuals = 0;
     CXXRecordDecl *ThisContext = nullptr;
     if (CXXMethodDecl *Method = dyn_cast<CXXMethodDecl>(Function)) {
       ThisContext = Method->getParent();
-      ThisTypeQuals = Method->getTypeQualifiers();
+      ThisTypeQuals = Method->getTypeQualifiers().getCVRQualifiers();;
     }
 
     CXXThisScopeRAII ThisScope(*this, ThisContext, ThisTypeQuals,
                                getLangOpts().CPlusPlus11);
 
     ResultType =
         SubstType(Proto->getReturnType(),
                   MultiLevelTemplateArgumentList(*ExplicitArgumentList),
AddImplicitObjectParameterType(ASTContext &Context,
   // C++11 [temp.func.order]p3:
   //   [...] The new parameter is of type "reference to cv A," where cv are
   //   the cv-qualifiers of the function template (if any) and A is
   //   the class of which the function template is a member.
   //
   // The standard doesn't say explicitly, but we pick the appropriate kind of
   // reference type based on [over.match.funcs]p4.
   QualType ArgTy = Context.getTypeDeclType(Method->getParent());
-  ArgTy = Context.getQualifiedType(ArgTy,
-                        Qualifiers::fromCVRMask(Method->getTypeQualifiers()));
+  ArgTy = Context.getQualifiedType(
+      ArgTy,
+      Qualifiers::fromCVRMask(Method->getTypeQualifiers().getCVRQualifiers()));
   if (Method->getRefQualifier() == RQ_RValue)
     ArgTy = Context.getRValueReferenceType(ArgTy);
   else
     ArgTy = Context.getLValueReferenceType(ArgTy);
   ArgTypes.push_back(ArgTy);
 }
 
 /// Determine whether the function template \p FT1 is at least as

lib/Sema/SemaTemplateInstantiateDecl.cpp

TemplateDeclInstantiator::SubstFunctionType(FunctionDecl *D,
   TypeSourceInfo *OldTInfo = D->getTypeSourceInfo();
   assert(OldTInfo && "substituting function without type source info");
   assert(Params.empty() && "parameter vector is non-empty at start");
 
   CXXRecordDecl *ThisContext = nullptr;
   unsigned ThisTypeQuals = 0;
   if (CXXMethodDecl *Method = dyn_cast<CXXMethodDecl>(D)) {
     ThisContext = cast<CXXRecordDecl>(Owner);
-    ThisTypeQuals = Method->getTypeQualifiers();
+    ThisTypeQuals = Method->getTypeQualifiers().getCVRQualifiers();
   }
 
   TypeSourceInfo *NewTInfo
     = SemaRef.SubstFunctionDeclType(OldTInfo, TemplateArgs,
                                     D->getTypeSpecStartLoc(),
                                     D->getDeclName(),
                                     ThisContext, ThisTypeQuals);
   if (!NewTInfo)

lib/Sema/SemaType.cpp

 
 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();
   ASTContext &Context = S.Context;
   const LangOptions &LangOpts = S.getLangOpts();
 
   // The name we're declaring, if any.
case DeclaratorChunk::Function: {
                                       FTI.getExceptionSpecType(),
                                       DynamicExceptions,
                                       DynamicExceptionRanges,
                                       NoexceptExpr,
                                       Exceptions,
                                       EPI.ExceptionSpec);
 
         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;
 
       // Handle pointer nullability.
       inferPointerNullability(SimplePointerKind::MemberPointer, DeclType.Loc,

 }
 
 /// GetTypeForDeclarator - Convert the type for the specified
 /// declarator to Type instances.
 ///
 /// The result of this call will never be null, but the associated
 /// type may be a null type if there's an unrecoverable error.
 TypeSourceInfo *Sema::GetTypeForDeclarator(Declarator &D, Scope *S) {
   // Determine the type of the declarator. Not all forms of declarator
   // have a type.
 
   TypeProcessingState state(*this, D);
 
   TypeSourceInfo *ReturnTypeInfo = nullptr;
   QualType T = GetDeclSpecTypeForDeclarator(state, ReturnTypeInfo);
   if (D.isPrototypeContext() && getLangOpts().ObjCAutoRefCount)
     inferARCWriteback(state, T);
bool IsPointee =
        D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::BlockPointer ||
        D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Reference);
   bool IsFuncReturnType =
       ChunkIndex > 0 &&
       D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Function;
   bool IsFuncType =
       ChunkIndex < D.getNumTypeObjects() &&
       D.getTypeObject(ChunkIndex).Kind == DeclaratorChunk::Function;
+  bool IsClassMemberFn = IsFuncType && D.getContext() == DeclaratorContext::MemberContext;
+
   if ( // Do not deduce addr space for function return type and function type,
-       // otherwise it will fail some sema check.
-      IsFuncReturnType || IsFuncType ||
+       // otherwise it will fail some sema check. We want to deduce class member functions.
+      IsFuncReturnType ||
+      (IsFuncType && !IsClassMemberFn) ||
       // Do not deduce addr space for member types of struct, except the pointee
-      // type of a pointer member type.
-      (D.getContext() == DeclaratorContext::MemberContext && !IsPointee) ||
+      // type of a pointer member type. We want to deduce class member functions.
+      (D.getContext() == DeclaratorContext::MemberContext && !IsPointee &&
+       !IsClassMemberFn) ||
       // 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))
if (State.getSema().getLangOpts().OpenCLVersion <= 120 &&
     // address spaces for some cases:
     // OpenCL 2.0, section 6.5:
     // The address space for a variable at program scope or a static variable
     // inside a function can either be __global or __constant, but defaults to
     // __global if not specified.
     // (...)
     // Pointers that are declared without pointing to a named address space
     // point to the generic address space.
-    if (IsPointee) {
+    // Deduce class members functions to be of the generic address space
+    if (IsPointee || IsClassMemberFn) {
       ImpAddr = LangAS::opencl_generic;
     } else {
       if (D.getContext() == DeclaratorContext::FileContext) {
         ImpAddr = LangAS::opencl_global;
       } else {
         if (D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_static ||
             D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_extern) {
           ImpAddr = LangAS::opencl_global;

test/CodeGenOpenCLCXX/addrspace-of-this.cl

+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s
+// expected-no-diagnostics
+
+// Test that the 'this' pointer is in the __generic address space.
+
+// FIXME: Add support for __constant address space.
+
+class C {
+public:
+  int v;
+  C() { v = 2; }
+  C(const C &c) { v = c.v; }
+  C &operator=(const C &c) {
+    v = c.v;
+    return *this;
+  }
+  int get() { return v; }
+};
+
+__global C c;
+
+__kernel void test__global() {
+  int i = c.get();
+  C c1(c);
+  C c2;
+  c2 = c1;
+}
+
+// CHECK-LABEL: @__cxx_global_var_init()
+// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) #4
+
+// Test that the address space is __generic for the constructor
+// CHECK-LABEL: @_ZN1CC1Ev(%class.C addrspace(4)* %this)
+// CHECK: entry:
+// CHECK:   %this.addr = alloca %class.C addrspace(4)*, align 4
+// CHECK:   store %class.C addrspace(4)* %this, %class.C addrspace(4)** %this.addr, align 4
+// CHECK:   %this1 = load %class.C addrspace(4)*, %class.C addrspace(4)** %this.addr, align 4
+// CHECK:   call void @_ZN1CC2Ev(%class.C addrspace(4)* %this1) #4
+// CHECK:   ret void
+
+// CHECK-LABEL: @_Z12test__globalv()
+
+// Test the address space of 'this' when invoking a method.
+// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+
+// Test the address space of 'this' when invoking copy-constructor.
+// CHECK: %0 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %0, %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+
+// Test the address space of 'this' when invoking a constructor.
+// CHECK:   %1 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK:   call void @_ZN1CC1Ev(%class.C addrspace(4)* %1) #4
+
+// Test the address space of 'this' when invoking assignment operator.
+// CHECK:   %2 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK:   %3 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK:   %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZN1CaSERU3AS4KS_(%class.C addrspace(4)* %3, %class.C addrspace(4)* dereferenceable(4) %2)
+
+#define TEST(AS)             \
+  __kernel void test##AS() { \
+    AS C c;                  \
+    int i = c.get();         \
+    C c1(c);                 \
+    C c2;                    \
+    c2 = c1;                 \
+  }
+
+TEST(__local)
+
+// CHECK-LABEL: _Z11test__localv
+// CHECK: @__cxa_guard_acquire
+
+// Test the address space of 'this' when invoking a method.
+// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
+
+// Test the address space of 'this' when invoking copy-constructor.
+// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
+
+// Test the address space of 'this' when invoking a constructor.
+// CHECK: %3 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %3)
+
+// Test the address space of 'this' when invoking assignment operator.
+// CHECK:  %4 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK:  %5 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK:  %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZN1CaSERU3AS4KS_(%class.C addrspace(4)* %5, %class.C addrspace(4)* dereferenceable(4) %4)
+
+TEST(__private)
+
+// CHECK-LABEL: @_Z13test__privatev
+
+// Test the address space of 'this' when invoking a method.
+// CHECK:   %1 = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK:   %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* %1)
+
+// Test the address space of 'this' when invoking a copy-constructor.
+// CHECK: %2 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: %3 = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %2, %class.C addrspace(4)* dereferenceable(4) %3)
+
+// Test the address space of 'this' when invoking a constructor.
+// CHECK: %4 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK:   call void @_ZN1CC1Ev(%class.C addrspace(4)* %4)
+
+// Test the address space of 'this' when invoking a copy-assignment.
+// CHECK:   %5 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK:   %6 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK:   %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZN1CaSERU3AS4KS_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %5)
+
+TEST()
+
+// CHECK-LABEL: @_Z4testv()
+// Test the address space of 'this' when invoking a method.
+// CHECK: %1 = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: %call = call i32 @_ZN1C3getEv(%class.C addrspace(4)* %1) #4
+
+// Test the address space of 'this' when invoking a copy-constructor.
+// CHECK: %2 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: %3 = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1ERU3AS4KS_(%class.C addrspace(4)* %2, %class.C addrspace(4)* dereferenceable(4) %3)
+
+// Test the address space of 'this' when invoking a constructor.
+// CHECK: %4 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZN1CC1Ev(%class.C addrspace(4)* %4)
+
+// Test the address space of 'this' when invoking a copy-assignment.
+// CHECK: %5 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: %6 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZN1CaSERU3AS4KS_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %5)

tools/libclang/CIndex.cpp

 
 unsigned clang_CXXMethod_isConst(CXCursor C) {
   if (!clang_isDeclaration(C.kind))
     return 0;
 
   const Decl *D = cxcursor::getCursorDecl(C);
   const CXXMethodDecl *Method =
       D ? dyn_cast_or_null<CXXMethodDecl>(D->getAsFunction()) : nullptr;
-  return (Method && (Method->getTypeQualifiers() & Qualifiers::Const)) ? 1 : 0;
+  return (Method && Method->getTypeQualifiers().hasConst()) ? 1 : 0;
 }
 
 unsigned clang_CXXMethod_isDefaulted(CXCursor C) {
   if (!clang_isDeclaration(C.kind))
     return 0;
 
   const Decl *D = cxcursor::getCursorDecl(C);
   const CXXMethodDecl *Method =