[OpenCL] Add generic AS to 'this' pointer

[mikael]

Summary

Address spaces are cast into generic before invoking the constructor.

[rjmccall]

lib/CodeGen/CGCall.cpp

81

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.

4052

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

81

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

81

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

4052

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

[rjmccall]

lib/CodeGen/CGCall.cpp

4052

Sorry, performAddrSpaceCast.

[Anastasia]

lib/CodeGen/CGCall.cpp

81

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

2191–2192

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

72–73

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

5261

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

[rjmccall]

include/clang/AST/DeclCXX.h

2193

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

4829

Can this be moved into deduceOpenCLImplicitAddrSpace?

[mikael]

lib/Sema/SemaType.cpp

4829

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

4829

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

209

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

803–804

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

2619–2620

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

275

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

1031

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.

3740

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

12000

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.

12367

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.

5071

You should be using the full qualifiers here.

lib/Sema/SemaTemplateDeduction.cpp

4662

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

275

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

2959

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

2621

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

275

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

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

lib/Sema/SemaExprCXX.cpp

1112

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

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

[mikael]

lib/AST/Type.cpp

2959

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

2621

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

2959

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

2959

Great, looks good.

lib/CodeGen/CGDebugInfo.cpp

2621

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

275

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

4830

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–4279

Another place where the address space was removed.

test/SemaOpenCLCXX/address-space-templates.cl

7–9

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

3931

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

5170

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

lib/Sema/TreeTransform.h

4277–4279

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–9

Do you know why it appears twice?

[mikael]

test/SemaOpenCLCXX/address-space-templates.cl

7–9

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

5170

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

5170

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

275

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

77

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

[rjmccall]

lib/CodeGen/CGCall.cpp

77

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/CanonicalType.h

struct CanProxyAdaptor<FunctionProtoType>
   LLVM_CLANG_CANPROXY_SIMPLE_ACCESSOR(
             ArrayRef<FunctionProtoType::ExtParameterInfo>, getExtParameterInfos)
 
   CanQualType getParamType(unsigned i) const {
     return CanQualType::CreateUnsafe(this->getTypePtr()->getParamType(i));
   }
 
   LLVM_CLANG_CANPROXY_SIMPLE_ACCESSOR(bool, isVariadic)
-  LLVM_CLANG_CANPROXY_SIMPLE_ACCESSOR(unsigned, getTypeQuals)
+  LLVM_CLANG_CANPROXY_SIMPLE_ACCESSOR(Qualifiers, getTypeQuals)
 
   using param_type_iterator =
       CanTypeIterator<FunctionProtoType::param_type_iterator>;
 
   param_type_iterator param_type_begin() const {
     return param_type_iterator(this->getTypePtr()->param_type_begin());
   }
 

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 {
+  static QualType getThisType(const FunctionProtoType *FPT,
+                              const CXXRecordDecl *Decl);
+
+  Qualifiers getTypeQualifiers() const {
     return getType()->getAs<FunctionProtoType>()->getTypeQuals();
   }
 
   /// 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 {
   ///   void f() &;
   ///   void g() &&;
   ///   void h();
   /// };

include/clang/AST/Type.h

public:
   }
 
   // Serialize these qualifiers into an opaque representation.
   unsigned getAsOpaqueValue() const {
     return Mask;
   }
 
   bool hasConst() const { return Mask & Const; }
-  void setConst(bool flag) {
-    Mask = (Mask & ~Const) | (flag ? Const : 0);
-  }
+  bool hasOnlyConst() const { return Mask == Const; }
   void removeConst() { Mask &= ~Const; }
   void addConst() { Mask |= Const; }
 
   bool hasVolatile() const { return Mask & Volatile; }
-  void setVolatile(bool flag) {
-    Mask = (Mask & ~Volatile) | (flag ? Volatile : 0);
-  }
+  bool hasOnlyVolatile() const { return Mask == Volatile; }
   void removeVolatile() { Mask &= ~Volatile; }
   void addVolatile() { Mask |= Volatile; }
 
   bool hasRestrict() const { return Mask & Restrict; }
-  void setRestrict(bool flag) {
-    Mask = (Mask & ~Restrict) | (flag ? Restrict : 0);
-  }
+  bool hasOnlyRestrict() const { return Mask == Restrict; }
   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;
   }
class FunctionTypeBitfields {
     unsigned RefQualifier : 2;
 
     /// Used only by FunctionProtoType, put here to pack with the
     /// other bitfields.
     /// The qualifiers are part of FunctionProtoType because...
     ///
     /// C++ 8.3.5p4: The return type, the parameter type list and the
     /// cv-qualifier-seq, [...], are part of the function type.
-    unsigned TypeQuals : 4;
+    unsigned FastTypeQuals : Qualifiers::FastWidth;
+    /// Whether this function has extended Qualifiers.
+    unsigned HasExtQuals : 1;
 
     /// The number of parameters this function has, not counting '...'.
     /// According to [implimits] 8 bits should be enough here but this is
     /// somewhat easy to exceed with metaprogramming and so we would like to
     /// keep NumParams as wide as reasonably possible.
     unsigned NumParams : 16;
 
     /// The type of exception specification this function has.
FunctionType(TypeClass tc, QualType res,
                bool VariablyModified, bool ContainsUnexpandedParameterPack,
                ExtInfo Info)
       : Type(tc, Canonical, Dependent, InstantiationDependent, VariablyModified,
              ContainsUnexpandedParameterPack),
         ResultType(res) {
     FunctionTypeBits.ExtInfo = Info.Bits;
   }
 
-  unsigned getTypeQuals() const { return FunctionTypeBits.TypeQuals; }
+  Qualifiers getFastTypeQuals() const {
+    return Qualifiers::fromFastMask(FunctionTypeBits.FastTypeQuals);
+  }
 
 public:
   QualType getReturnType() const { return ResultType; }
 
   bool getHasRegParm() const { return getExtInfo().getHasRegParm(); }
   unsigned getRegParmType() const { return getExtInfo().getRegParm(); }
 
   /// Determine whether this function type includes the GNU noreturn
   /// attribute. The C++11 [[noreturn]] attribute does not affect the function
   /// type.
   bool getNoReturnAttr() const { return getExtInfo().getNoReturn(); }
 
   CallingConv getCallConv() const { return getExtInfo().getCC(); }
   ExtInfo getExtInfo() const { return ExtInfo(FunctionTypeBits.ExtInfo); }
-  bool isConst() const { return getTypeQuals() & Qualifiers::Const; }
-  bool isVolatile() const { return getTypeQuals() & Qualifiers::Volatile; }
-  bool isRestrict() const { return getTypeQuals() & Qualifiers::Restrict; }
+
+  static_assert((~Qualifiers::FastMask & Qualifiers::CVRMask) == 0,
+                "Const, volatile and restrict are assumed to be a subset of "
+                "the fast qualifiers.");
+
+  bool isConst() const { return getFastTypeQuals().hasConst(); }
+  bool isVolatile() const { return getFastTypeQuals().hasVolatile(); }
+  bool isRestrict() const { return getFastTypeQuals().hasRestrict(); }
 
   /// Determine the type of an expression that calls a function of
   /// this type.
   QualType getCallResultType(const ASTContext &Context) const {
     return getReturnType().getNonLValueExprType(Context);
   }
 
   static StringRef getNameForCallConv(CallingConv CC);

 /// 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> {
+          FunctionType::ExtParameterInfo, Qualifiers> {
   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,
class FunctionProtoType final
   //   * a single FunctionDecl *
   //   used to store information about the various types of exception
   //   specification. See getExceptionSpecSize for the details.
   //
   // * Optionally an array of getNumParams() ExtParameterInfo holding
   //   an ExtParameterInfo for each of the parameters. Present if and
   //   only if hasExtParameterInfos() is true.
   //
+  // * Optionally a Qualifiers object to represent extra qualifiers that can't
+  //   be represented by FunctionTypeBitfields.FastTypeQuals. Present if and only
+  //   if hasExtQualifiers() is true.
+  //
   // The optional FunctionTypeExtraBitfields has to be before the data
   // related to the exception specification since it contains the number
   // of exception types.
   //
   // We put the ExtParameterInfos last.  If all were equal, it would make
   // more sense to put these before the exception specification, because
   // it's much easier to skip past them compared to the elaborate switch
   // required to skip the exception specification.  However, all is not
public:
 
   /// Extra information about a function prototype. ExtProtoInfo is not
   /// stored as such in FunctionProtoType but is used to group together
   /// the various bits of extra information about a function prototype.
   struct ExtProtoInfo {
     FunctionType::ExtInfo ExtInfo;
     bool Variadic : 1;
     bool HasTrailingReturn : 1;
-    unsigned char TypeQuals = 0;
+    Qualifiers TypeQuals;
     RefQualifierKind RefQualifier = RQ_None;
     ExceptionSpecInfo ExceptionSpec;
     const ExtParameterInfo *ExtParameterInfos = nullptr;
 
     ExtProtoInfo() : Variadic(false), HasTrailingReturn(false) {}
 
     ExtProtoInfo(CallingConv CC)
         : ExtInfo(CC), Variadic(false), HasTrailingReturn(false) {}
static bool hasExtraBitfields(ExceptionSpecificationType EST) {
     return EST == EST_Dynamic;
   }
 
   /// Whether the trailing FunctionTypeExtraBitfields is present.
   bool hasExtraBitfields() const {
     return hasExtraBitfields(getExceptionSpecType());
   }
 
+  bool hasExtQualifiers() const {
+    return FunctionTypeBits.HasExtQuals;
+  }
+
 public:
   unsigned getNumParams() const { return FunctionTypeBits.NumParams; }
 
   QualType getParamType(unsigned i) const {
     assert(i < getNumParams() && "invalid parameter index");
     return param_type_begin()[i];
   }
 
   ArrayRef<QualType> getParamTypes() const {
     return llvm::makeArrayRef(param_type_begin(), param_type_end());
   }
 
   ExtProtoInfo getExtProtoInfo() const {
     ExtProtoInfo EPI;
     EPI.ExtInfo = getExtInfo();
     EPI.Variadic = isVariadic();
     EPI.HasTrailingReturn = hasTrailingReturn();
     EPI.ExceptionSpec.Type = getExceptionSpecType();
-    EPI.TypeQuals = static_cast<unsigned char>(getTypeQuals());
+    EPI.TypeQuals = getTypeQuals();
     EPI.RefQualifier = getRefQualifier();
     if (EPI.ExceptionSpec.Type == EST_Dynamic) {
       EPI.ExceptionSpec.Exceptions = exceptions();
     } else if (isComputedNoexcept(EPI.ExceptionSpec.Type)) {
       EPI.ExceptionSpec.NoexceptExpr = getNoexceptExpr();
     } else if (EPI.ExceptionSpec.Type == EST_Uninstantiated) {
       EPI.ExceptionSpec.SourceDecl = getExceptionSpecDecl();
       EPI.ExceptionSpec.SourceTemplate = getExceptionSpecTemplate();
public:
   /// A function template whose last parameter is a parameter pack can be
   /// called with an arbitrary number of arguments, much like a variadic
   /// function.
   bool isTemplateVariadic() const;
 
   /// Whether this function prototype has a trailing return type.
   bool hasTrailingReturn() const { return FunctionTypeBits.HasTrailingReturn; }
 
-  unsigned getTypeQuals() const { return FunctionType::getTypeQuals(); }
+  Qualifiers getTypeQuals() const {
+    if (hasExtQualifiers())
+      return *getTrailingObjects<Qualifiers>();
+    else
+      return getFastTypeQuals();
+  }
 
   /// Retrieve the ref-qualifier associated with this function type.
   RefQualifierKind getRefQualifier() const {
     return static_cast<RefQualifierKind>(FunctionTypeBits.RefQualifier);
   }
 
   using param_type_iterator = const QualType *;
   using param_type_range = llvm::iterator_range<param_type_iterator>;

include/clang/Sema/Sema.h

class CXXThisScopeRAII {
     QualType OldCXXThisTypeOverride;
     bool Enabled;
 
   public:
     /// Introduce a new scope where 'this' may be allowed (when enabled),
     /// using the given declaration (which is either a class template or a
     /// class) along with the given qualifiers.
     /// along with the qualifiers placed on '*this'.
-    CXXThisScopeRAII(Sema &S, Decl *ContextDecl, unsigned CXXThisTypeQuals,
+    CXXThisScopeRAII(Sema &S, Decl *ContextDecl, Qualifiers CXXThisTypeQuals,
                      bool Enabled = true);
 
     ~CXXThisScopeRAII();
   };
 
   /// Make sure the value of 'this' is actually available in the current
   /// context, if it is a potentially evaluated context.
   ///
TypeSourceInfo *SubstType(TypeLoc TL,
                             const MultiLevelTemplateArgumentList &TemplateArgs,
                             SourceLocation Loc, DeclarationName Entity);
 
   TypeSourceInfo *SubstFunctionDeclType(TypeSourceInfo *T,
                             const MultiLevelTemplateArgumentList &TemplateArgs,
                                         SourceLocation Loc,
                                         DeclarationName Entity,
                                         CXXRecordDecl *ThisContext,
-                                        unsigned ThisTypeQuals);
+                                        Qualifiers ThisTypeQuals);
   void SubstExceptionSpec(FunctionDecl *New, const FunctionProtoType *Proto,
                           const MultiLevelTemplateArgumentList &Args);
   bool SubstExceptionSpec(SourceLocation Loc,
                           FunctionProtoType::ExceptionSpecInfo &ESI,
                           SmallVectorImpl<QualType> &ExceptionStorage,
                           const MultiLevelTemplateArgumentList &Args);
   ParmVarDecl *SubstParmVarDecl(ParmVarDecl *D,
                             const MultiLevelTemplateArgumentList &TemplateArgs,

lib/AST/ASTContext.cpp

QualType ASTContext::getFunctionTypeInternal(
 
   // Compute the needed size to hold this FunctionProtoType and the
   // various trailing objects.
   auto ESH = FunctionProtoType::getExceptionSpecSize(
       EPI.ExceptionSpec.Type, EPI.ExceptionSpec.Exceptions.size());
   size_t Size = FunctionProtoType::totalSizeToAlloc<
       QualType, FunctionType::FunctionTypeExtraBitfields,
       FunctionType::ExceptionType, Expr *, FunctionDecl *,
-      FunctionProtoType::ExtParameterInfo>(
+      FunctionProtoType::ExtParameterInfo, Qualifiers>(
       NumArgs, FunctionProtoType::hasExtraBitfields(EPI.ExceptionSpec.Type),
       ESH.NumExceptionType, ESH.NumExprPtr, ESH.NumFunctionDeclPtr,
-      EPI.ExtParameterInfos ? NumArgs : 0);
+      EPI.ExtParameterInfos ? NumArgs : 0,
+      EPI.TypeQuals.hasNonFastQualifiers() ? 1 : 0);
 
   auto *FTP = (FunctionProtoType *)Allocate(Size, TypeAlignment);
   FunctionProtoType::ExtProtoInfo newEPI = EPI;
   new (FTP) FunctionProtoType(ResultTy, ArgArray, Canonical, newEPI);
   Types.push_back(FTP);
   if (!Unique)
     FunctionProtoTypes.InsertNode(FTP, InsertPos);
   return QualType(FTP, 0);

lib/AST/ASTDumper.cpp

void VisitFunctionType(const FunctionType *T) {
       if (EI.getProducesResult()) OS << " produces_result";
       if (EI.getHasRegParm()) OS << " regparm " << EI.getRegParm();
       OS << " " << FunctionType::getNameForCallConv(EI.getCC());
       dumpTypeAsChild(T->getReturnType());
     }
     void VisitFunctionProtoType(const FunctionProtoType *T) {
       auto EPI = T->getExtProtoInfo();
       if (EPI.HasTrailingReturn) OS << " trailing_return";
-      if (T->isConst()) OS << " const";
-      if (T->isVolatile()) OS << " volatile";
-      if (T->isRestrict()) OS << " restrict";
+
+      if (!T->getTypeQuals().empty())
+        OS << " " << T->getTypeQuals().getAsString();
+
       switch (EPI.RefQualifier) {
         case RQ_None: break;
         case RQ_LValue: OS << " &"; break;
         case RQ_RValue: OS << " &&"; break;
       }
       // FIXME: Exception specification.
       // FIXME: Consumed parameters.
       VisitFunctionType(T);
       for (QualType PT : T->getParamTypes())
         dumpTypeAsChild(PT);
       if (EPI.Variadic)
         dumpChild([=] { OS << "..."; });

lib/AST/DeclCXX.cpp

 
 CXXMethodDecl::overridden_method_range
 CXXMethodDecl::overridden_methods() const {
   if (isa<CXXConstructorDecl>(this))
     return overridden_method_range(nullptr, nullptr);
   return getASTContext().overridden_methods(this);
 }
 
+QualType CXXMethodDecl::getThisType(const FunctionProtoType *FPT,
+                                    const CXXRecordDecl *Decl) {
+  ASTContext &C = Decl->getASTContext();
+  QualType ClassTy = C.getTypeDeclType(Decl);
+  ClassTy = C.getQualifiedType(ClassTy, FPT->getTypeQuals());
+  return C.getPointerType(ClassTy);
+}
+
 QualType CXXMethodDecl::getThisType(ASTContext &C) const {
   // C++ 9.3.2p1: The type of this in a member function of a class X is X*.
   // 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()));
-  return C.getPointerType(ClassTy);
+  return CXXMethodDecl::getThisType(getType()->getAs<FunctionProtoType>(),
+                                    getParent());
 }
 
 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)
     CheckFn = this;

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 = Method->getTypeQualifiers();
     // 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.
   const TemplateArgumentList *TemplateArgs = nullptr;

 // <type>          ::= <function-type>
 // <function-type> ::= [<CV-qualifiers>] F [Y]
 //                      <bare-function-type> [<ref-qualifier>] E
 void CXXNameMangler::mangleType(const FunctionProtoType *T) {
   mangleExtFunctionInfo(T);
 
   // Mangle CV-qualifiers, if present.  These are 'this' qualifiers,
   // e.g. "const" in "int (A::*)() const".
-  mangleQualifiers(Qualifiers::fromCVRUMask(T->getTypeQuals()));
+  mangleQualifiers(T->getTypeQuals());
 
   // Mangle instantiation-dependent exception-specification, if present,
   // per cxx-abi-dev proposal on 2016-10-11.
   if (T->hasInstantiationDependentExceptionSpec()) {
     if (isComputedNoexcept(T->getExceptionSpecType())) {
       Out << "DO";
       mangleExpression(T->getNoexceptExpr());
       Out << "E";

lib/AST/MicrosoftMangle.cpp

if (isa<CXXDestructorDecl>(MD)) {
         CC = getASTContext().getDefaultCallingConvention(
             /*IsVariadic=*/false, /*IsCXXMethod=*/true);
     }
   }
 
   // If this is a C++ instance method, mangle the CVR qualifiers for the
   // this pointer.
   if (HasThisQuals) {
-    Qualifiers Quals = Qualifiers::fromCVRUMask(Proto->getTypeQuals());
+    Qualifiers Quals = Proto->getTypeQuals();
     manglePointerExtQualifiers(Quals, /*PointeeType=*/QualType());
     mangleRefQualifier(Proto->getRefQualifier());
     mangleQualifiers(Quals, /*IsMember=*/false);
   }
 
   mangleCallingConvention(CC);
 
   // <return-type> ::= <type>

lib/AST/Type.cpp

 
 FunctionProtoType::FunctionProtoType(QualType result, ArrayRef<QualType> params,
                                      QualType canonical,
                                      const ExtProtoInfo &epi)
     : FunctionType(FunctionProto, result, canonical, result->isDependentType(),
                    result->isInstantiationDependentType(),
                    result->isVariablyModifiedType(),
                    result->containsUnexpandedParameterPack(), epi.ExtInfo) {
-  FunctionTypeBits.TypeQuals = epi.TypeQuals;
+  FunctionTypeBits.FastTypeQuals = epi.TypeQuals.getFastQualifiers();
   FunctionTypeBits.RefQualifier = epi.RefQualifier;
   FunctionTypeBits.NumParams = params.size();
   assert(getNumParams() == params.size() && "NumParams overflow!");
   FunctionTypeBits.ExceptionSpecType = epi.ExceptionSpec.Type;
   FunctionTypeBits.HasExtParameterInfos = !!epi.ExtParameterInfos;
   FunctionTypeBits.Variadic = epi.Variadic;
   FunctionTypeBits.HasTrailingReturn = epi.HasTrailingReturn;
 
FunctionProtoType::FunctionProtoType(QualType result, ArrayRef<QualType> params,
   }
 
   // Fill in the extra parameter info if present.
   if (epi.ExtParameterInfos) {
     auto *extParamInfos = getTrailingObjects<ExtParameterInfo>();
     for (unsigned i = 0; i != getNumParams(); ++i)
       extParamInfos[i] = epi.ExtParameterInfos[i];
   }
+
+  if (epi.TypeQuals.hasNonFastQualifiers()) {
+    FunctionTypeBits.HasExtQuals = 1;
+    *getTrailingObjects<Qualifiers>() = epi.TypeQuals;
+  } else {
+    FunctionTypeBits.HasExtQuals = 0;
+  }
 }
 
 bool FunctionProtoType::hasDependentExceptionSpec() const {
   if (Expr *NE = getNoexceptExpr())
     return NE->isValueDependent();
   for (QualType ET : exceptions())
     // A pack expansion with a non-dependent pattern is still dependent,
     // because we don't know whether the pattern is in the exception spec
void FunctionProtoType::Profile(llvm::FoldingSetNodeID &ID, QualType Result,
 
   ID.AddPointer(Result.getAsOpaquePtr());
   for (unsigned i = 0; i != NumParams; ++i)
     ID.AddPointer(ArgTys[i].getAsOpaquePtr());
   // This method is relatively performance sensitive, so as a performance
   // shortcut, use one AddInteger call instead of four for the next four
   // fields.
   assert(!(unsigned(epi.Variadic) & ~1) &&
-         !(unsigned(epi.TypeQuals) & ~255) &&
          !(unsigned(epi.RefQualifier) & ~3) &&
          !(unsigned(epi.ExceptionSpec.Type) & ~15) &&
          "Values larger than expected.");
   ID.AddInteger(unsigned(epi.Variadic) +
-                (epi.TypeQuals << 1) +
-                (epi.RefQualifier << 9) +
-                (epi.ExceptionSpec.Type << 11));
+                (epi.RefQualifier << 1) +
+                (epi.ExceptionSpec.Type << 3));
+  ID.Add(epi.TypeQuals);
   if (epi.ExceptionSpec.Type == EST_Dynamic) {
     for (QualType Ex : epi.ExceptionSpec.Exceptions)
       ID.AddPointer(Ex.getAsOpaquePtr());
   } else if (isComputedNoexcept(epi.ExceptionSpec.Type)) {
     epi.ExceptionSpec.NoexceptExpr->Profile(ID, Context, Canonical);
   } else if (epi.ExceptionSpec.Type == EST_Uninstantiated ||
              epi.ExceptionSpec.Type == EST_Unevaluated) {
     ID.AddPointer(epi.ExceptionSpec.SourceDecl->getCanonicalDecl());

lib/AST/TypePrinter.cpp

if (T->isVariadic()) {
     OS << "void";
   }
 
   OS << ')';
 
   FunctionType::ExtInfo Info = T->getExtInfo();
 
   printFunctionAfter(Info, OS);
 
-  if (unsigned quals = T->getTypeQuals()) {
-    OS << ' ';
-    AppendTypeQualList(OS, quals, Policy.Restrict);
-  }
+  if (!T->getTypeQuals().empty())
+    OS << " " << T->getTypeQuals().getAsString();
 
   switch (T->getRefQualifier()) {
   case RQ_None:
     break;
 
   case RQ_LValue:
     OS << " &";
     break;

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/CodeGenOptions.h"
 #include "clang/Basic/TargetBuiltins.h"
 #include "clang/CodeGen/CGFunctionInfo.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/CGDebugInfo.cpp

llvm::DIType *CGDebugInfo::CreateType(const MemberPointerType *Ty,
   llvm::DIType *ClassType = getOrCreateType(QualType(Ty->getClass(), 0), U);
   if (Ty->isMemberDataPointerType())
     return DBuilder.createMemberPointerType(
         getOrCreateType(Ty->getPointeeType(), U), ClassType, Size, /*Align=*/0,
         Flags);
 
   const FunctionProtoType *FPT =
       Ty->getPointeeType()->getAs<FunctionProtoType>();
   return DBuilder.createMemberPointerType(
-      getOrCreateInstanceMethodType(CGM.getContext().getPointerType(QualType(
-                                        Ty->getClass(), FPT->getTypeQuals())),
-                                    FPT, U),
+      getOrCreateInstanceMethodType(
+          CXXMethodDecl::getThisType(FPT, Ty->getMostRecentCXXRecordDecl()),
       ClassType, Size, /*Align=*/0, Flags);
 }
 
 llvm::DIType *CGDebugInfo::CreateType(const AtomicType *Ty, llvm::DIFile *U) {
   auto *FromTy = getOrCreateType(Ty->getValueType(), U);
   return DBuilder.createQualifiedType(llvm::dwarf::DW_TAG_atomic_type, FromTy);
 }
 

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/CodeGen/CGExpr.cpp

if (rec->isUnion()) {
     // For structs, we GEP to the field that the record layout suggests.
     addr = emitAddrOfFieldStorage(*this, addr, field);
 
     // If this is a reference field, load the reference right now.
     if (FieldType->isReferenceType()) {
       LValue RefLVal = MakeAddrLValue(addr, FieldType, FieldBaseInfo,
                                       FieldTBAAInfo);
       if (RecordCVR & Qualifiers::Volatile)
-        RefLVal.getQuals().setVolatile(true);
+        RefLVal.getQuals().addVolatile();
       addr = EmitLoadOfReference(RefLVal, &FieldBaseInfo, &FieldTBAAInfo);
 
       // Qualifiers on the struct don't apply to the referencee.
       RecordCVR = 0;
       FieldType = FieldType->getPointeeType();
     }
   }
 

lib/CodeGen/CGValue.h

public:
 
   Qualifiers getQualifiers() const { return Quals; }
 
   bool isVolatile() const {
     return Quals.hasVolatile();
   }
 
   void setVolatile(bool flag) {
-    Quals.setVolatile(flag);
+    if (flag)
+      Quals.addVolatile();
+    else
+      Quals.removeVolatile();
   }
 
   Qualifiers::ObjCLifetime getObjCLifetime() const {
     return Quals.getObjCLifetime();
   }
 
   NeedsGCBarriers_t requiresGCollection() const {
     return NeedsGCBarriers_t(ObjCGCFlag);

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())
+    // FIXME: OpenCL: Need to consider address spaces
+    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(),
                                      getLangOpts().CPlusPlus11);
 
     // Parse the exception-specification.
     SourceRange SpecificationRange;
     SmallVector<ParsedType, 4> DynamicExceptions;
     SmallVector<SourceRange, 4> DynamicExceptionRanges;
     ExprResult NoexceptExpr;
     CachedTokens *ExceptionSpecTokens;
void Parser::ParseLexedMemberInitializers(ParsingClass &Class) {
 
   if (!Class.LateParsedDeclarations.empty()) {
     // C++11 [expr.prim.general]p4:
     //   Otherwise, if a member-declarator declares a non-static data member
     //  (9.2) of a class X, the expression this is a prvalue of type "pointer
     //  to X" within the optional brace-or-equal-initializer. It shall not
     //  appear elsewhere in the member-declarator.
     Sema::CXXThisScopeRAII ThisScope(Actions, Class.TagOrTemplate,
-                                     /*TypeQuals=*/(unsigned)0);
+                                     Qualifiers());
 
     for (size_t i = 0; i < Class.LateParsedDeclarations.size(); ++i) {
       Class.LateParsedDeclarations[i]->ParseLexedMemberInitializers();
     }
   }
 
   if (!AlreadyHasClassScope)
     Actions.ActOnFinishDelayedMemberDeclarations(getCurScope(),

lib/Parse/ParseDecl.cpp

void Parser::ParseLexedAttribute(LateParsedAttribute &LA,
   SourceLocation endLoc;
 
   if (LA.Decls.size() > 0) {
     Decl *D = LA.Decls[0];
     NamedDecl *ND  = dyn_cast<NamedDecl>(D);
     RecordDecl *RD = dyn_cast_or_null<RecordDecl>(D->getDeclContext());
 
     // Allow 'this' within late-parsed attributes.
-    Sema::CXXThisScopeRAII ThisScope(Actions, RD, /*TypeQuals=*/0,
+    Sema::CXXThisScopeRAII ThisScope(Actions, RD, Qualifiers(),
                                      ND && ND->isCXXInstanceMember());
 
     if (LA.Decls.size() == 1) {
       // If the Decl is templatized, add template parameters to scope.
       bool HasTemplateScope = EnterScope && D->isTemplateDecl();
       ParseScope TempScope(this, Scope::TemplateParamScope, HasTemplateScope);
       if (HasTemplateScope)
         Actions.ActOnReenterTemplateScope(Actions.CurScope, D);
if (getLangOpts().CPlusPlus) {
       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());
-      Sema::CXXThisScopeRAII ThisScope(Actions,
-                               dyn_cast<CXXRecordDecl>(Actions.CurContext),
-                               DS.getTypeQualifiers() |
-                               (D.getDeclSpec().isConstexprSpecified() &&
-                                !getLangOpts().CPlusPlus14
-                                  ? Qualifiers::Const : 0),
-                               IsCXX11MemberFunction);
+
+      Qualifiers Q = Qualifiers::fromCVRUMask(DS.getTypeQualifiers());
+      if (D.getDeclSpec().isConstexprSpecified() && !getLangOpts().CPlusPlus14)
+        Q.addConst();
+
+      Sema::CXXThisScopeRAII ThisScope(
+          Actions, dyn_cast<CXXRecordDecl>(Actions.CurContext), Q,
 
       // Parse exception-specification[opt].
       bool Delayed = D.isFirstDeclarationOfMember() &&
                      D.isFunctionDeclaratorAFunctionDeclaration();
       if (Delayed && Actions.isLibstdcxxEagerExceptionSpecHack(D) &&
           GetLookAheadToken(0).is(tok::kw_noexcept) &&
           GetLookAheadToken(1).is(tok::l_paren) &&
           GetLookAheadToken(2).is(tok::kw_noexcept) &&

lib/Parse/ParseOpenMP.cpp

 public:
   FNContextRAII(Parser &P, Parser::DeclGroupPtrTy Ptr) : P(P) {
     Decl *D = *Ptr.get().begin();
     NamedDecl *ND = dyn_cast<NamedDecl>(D);
     RecordDecl *RD = dyn_cast_or_null<RecordDecl>(D->getDeclContext());
     Sema &Actions = P.getActions();
 
     // Allow 'this' within late-parsed attributes.
-    ThisScope = new Sema::CXXThisScopeRAII(Actions, RD, /*TypeQuals=*/0,
+    ThisScope = new Sema::CXXThisScopeRAII(Actions, RD, Qualifiers(),
                                            ND && ND->isCXXInstanceMember());
 
     // If the Decl is templatized, add template parameters to scope.
     HasTemplateScope = D->isTemplateDecl();
     TempScope =
         new Parser::ParseScope(&P, Scope::TemplateParamScope, HasTemplateScope);
     if (HasTemplateScope)
       Actions.ActOnReenterTemplateScope(Actions.getCurScope(), D);

lib/Sema/SemaCodeComplete.cpp

void ResultBuilder::AddResult(Result R, DeclContext *CurContext,
   if (InBaseClass)
     setInBaseClass(R);
 
   AdjustResultPriorityForDecl(R);
 
   if (HasObjectTypeQualifiers)
     if (const auto *Method = dyn_cast<CXXMethodDecl>(R.Declaration))
       if (Method->isInstance()) {
-        Qualifiers MethodQuals =
-            Qualifiers::fromCVRMask(Method->getTypeQualifiers());
+        Qualifiers MethodQuals = Method->getTypeQualifiers();
         if (ObjectTypeQualifiers == MethodQuals)
           R.Priority += CCD_ObjectQualifierMatch;
         else if (ObjectTypeQualifiers - MethodQuals) {
           // The method cannot be invoked, because doing so would drop
           // qualifiers.
           return;
         }
       }
AddFunctionTypeQualsToCompletionString(CodeCompletionBuilder &Result,
                                        const FunctionDecl *Function) {
   const auto *Proto = Function->getType()->getAs<FunctionProtoType>();
   if (!Proto || !Proto->getTypeQuals())
     return;
 
   // FIXME: Add ref-qualifier!
 
   // Handle single qualifiers without copying
-  if (Proto->getTypeQuals() == Qualifiers::Const) {
+  if (Proto->getTypeQuals().hasOnlyConst()) {
     Result.AddInformativeChunk(" const");
     return;
   }
 
-  if (Proto->getTypeQuals() == Qualifiers::Volatile) {
+  if (Proto->getTypeQuals().hasOnlyVolatile()) {
     Result.AddInformativeChunk(" volatile");
     return;
   }
 
-  if (Proto->getTypeQuals() == Qualifiers::Restrict) {
+  if (Proto->getTypeQuals().hasOnlyRestrict()) {
     Result.AddInformativeChunk(" restrict");
     return;
   }
 
   // Handle multiple qualifiers.
   std::string QualsStr;
   if (Proto->isConst())
     QualsStr += " const";
case PCC_RecoveryInFunction:
     // Unfiltered
     break;
   }
 
   // If we are in a C++ non-static member function, check the qualifiers on
   // the member function to filter/prioritize the results list.
   if (CXXMethodDecl *CurMethod = dyn_cast<CXXMethodDecl>(CurContext)) {
     if (CurMethod->isInstance()) {
-      Results.setObjectTypeQualifiers(
-          Qualifiers::fromCVRMask(CurMethod->getTypeQualifiers()));
+      Results.setObjectTypeQualifiers(CurMethod->getTypeQualifiers());
     }
   }
 
   CodeCompletionDeclConsumer Consumer(Results, CurContext);
   LookupVisibleDecls(S, LookupOrdinaryName, Consumer,
                      CodeCompleter->includeGlobals(),
                      CodeCompleter->loadExternal());
 

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));
+
+    QualType AdjustedQT = QualType(AdjustedType, 0);
+    LangAS AS = Old->getType().getAddressSpace();
+    AdjustedQT = Context.getAddrSpaceQualType(AdjustedQT, AS);
+
+    New->setType(AdjustedQT);
     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>() &&
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;
+      EPI.TypeQuals.addConst();
       MD->setType(Context.getFunctionType(FPT->getReturnType(),
                                           FPT->getParamTypes(), EPI));
 
       // Warn that we did this, if we're not performing template instantiation.
       // In that case, we'll have warned already when the template was defined.
       if (!inTemplateInstantiation()) {
         SourceLocation AddConstLoc;
         if (FunctionTypeLoc FTL = MD->getTypeSourceInfo()->getTypeLoc()

lib/Sema/SemaDeclCXX.cpp

QualType Sema::CheckConstructorDeclarator(Declarator &D, QualType R,
   // Rebuild the function type "R" without any type qualifiers (in
   // case any of the errors above fired) and with "void" as the
   // return type, since constructors don't have return types.
   const FunctionProtoType *Proto = R->getAs<FunctionProtoType>();
   if (Proto->getReturnType() == Context.VoidTy && !D.isInvalidType())
     return R;
 
   FunctionProtoType::ExtProtoInfo EPI = Proto->getExtProtoInfo();
-  EPI.TypeQuals = 0;
+  EPI.TypeQuals = Qualifiers();
   EPI.RefQualifier = RQ_None;
 
   return Context.getFunctionType(Context.VoidTy, Proto->getParamTypes(), EPI);
 }
 
 /// CheckConstructor - Checks a fully-formed constructor for
 /// well-formedness, issuing any diagnostics required. Returns true if
 /// the constructor declarator is invalid.
QualType Sema::CheckDestructorDeclarator(Declarator &D, QualType R,
   // "void" as the return type, since destructors don't have return
   // types.
   if (!D.isInvalidType())
     return R;
 
   const FunctionProtoType *Proto = R->getAs<FunctionProtoType>();
   FunctionProtoType::ExtProtoInfo EPI = Proto->getExtProtoInfo();
   EPI.Variadic = false;
-  EPI.TypeQuals = 0;
+  EPI.TypeQuals = Qualifiers();
   EPI.RefQualifier = RQ_None;
   return Context.getFunctionType(Context.VoidTy, None, EPI);
 }
 
 static void extendLeft(SourceRange &R, SourceRange Before) {
   if (Before.isInvalid())
     return;
   R.setBegin(Before.getBegin());
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(
+                   Context.getQualifiedType(
                        BaseType, CopyAssignOperator->getTypeQualifiers()),
                    VK_LValue, BasePath);
 
     // Build the copy.
     StmtResult Copy = buildSingleCopyAssign(*this, Loc, BaseType,
                                             To, From,
                                             /*CopyingBaseSubobject=*/true,
                                             /*Copying=*/true);
     if (Copy.isInvalid()) {
       CopyAssignOperator->setInvalidDecl();
       return;
     }
 
     // Success! Record the copy.
for (auto &Base : ClassDecl->bases()) {
     // 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(
+                   Context.getQualifiedType(
                        BaseType, MoveAssignOperator->getTypeQualifiers()),
                    VK_LValue, BasePath);
 
     // Build the move.
     StmtResult Move = buildSingleCopyAssign(*this, Loc, BaseType,
                                             To, From,
                                             /*CopyingBaseSubobject=*/true,
                                             /*Copying=*/false);
     if (Move.isInvalid()) {
       MoveAssignOperator->setInvalidDecl();
       return;
     }
 
     // Success! Record the move.

lib/Sema/SemaExpr.cpp

void Sema::ActOnBlockArguments(SourceLocation CaretLoc, Declarator &ParamInfo,
   QualType T = Sig->getType();
 
   // FIXME: We should allow unexpanded parameter packs here, but that would,
   // in turn, make the block expression contain unexpanded parameter packs.
   if (DiagnoseUnexpandedParameterPack(CaretLoc, Sig, UPPC_Block)) {
     // Drop the parameters.
     FunctionProtoType::ExtProtoInfo EPI;
     EPI.HasTrailingReturn = false;
-    EPI.TypeQuals |= DeclSpec::TQ_const;
+    EPI.TypeQuals.addConst();
     T = Context.getFunctionType(Context.DependentTy, None, EPI);
     Sig = Context.getTrivialTypeSourceInfo(T);
   }
 
   // GetTypeForDeclarator always produces a function type for a block
   // literal signature.  Furthermore, it is always a FunctionProtoType
   // unless the function was written with a typedef.
   assert(T->isFunctionType() &&
if (!BSI->FunctionType.isNull()) {
     } else if (FTy->getReturnType() == RetTy &&
                (!NoReturn || FTy->getNoReturnAttr())) {
       BlockTy = BSI->FunctionType;
 
     // Otherwise, make the minimal modifications to the function type.
     } else {
       const FunctionProtoType *FPT = cast<FunctionProtoType>(FTy);
       FunctionProtoType::ExtProtoInfo EPI = FPT->getExtProtoInfo();
-      EPI.TypeQuals = 0; // FIXME: silently?
+      EPI.TypeQuals = Qualifiers();
       EPI.ExtInfo = Ext;
       BlockTy = Context.getFunctionType(RetTy, FPT->getParamTypes(), EPI);
     }
 
   // If we don't have a function type, just build one from nothing.
   } else {
     FunctionProtoType::ExtProtoInfo EPI;
     EPI.ExtInfo = FunctionType::ExtInfo().withNoReturn(NoReturn);

lib/Sema/SemaExprCXX.cpp

QualType Sema::getCurrentThisType() {
   if (!ThisTy.isNull() && isLambdaCallOperator(CurContext))
     return adjustCVQualifiersForCXXThisWithinLambda(FunctionScopes, ThisTy,
                                                     CurContext, Context);
   return ThisTy;
 }
 
 Sema::CXXThisScopeRAII::CXXThisScopeRAII(Sema &S,
                                          Decl *ContextDecl,
-                                         unsigned CXXThisTypeQuals,
+                                         Qualifiers CXXThisTypeQuals,
                                          bool Enabled)
   : S(S), OldCXXThisTypeOverride(S.CXXThisTypeOverride), Enabled(false)
 {
   if (!Enabled || !ContextDecl)
     return;
 
   CXXRecordDecl *Record = nullptr;
   if (ClassTemplateDecl *Template = dyn_cast<ClassTemplateDecl>(ContextDecl))
     Record = Template->getTemplatedDecl();
   else
     Record = cast<CXXRecordDecl>(ContextDecl);
 
-  // We care only for CVR qualifiers here, so cut everything else.
-  CXXThisTypeQuals &= Qualifiers::FastMask;
-  S.CXXThisTypeOverride
-    = S.Context.getPointerType(
-        S.Context.getRecordType(Record).withCVRQualifiers(CXXThisTypeQuals));
+  QualType T = S.Context.getRecordType(Record);
+  T = S.getASTContext().getQualifiedType(T, CXXThisTypeQuals);
+
+  S.CXXThisTypeOverride = S.Context.getPointerType(T);
 
   this->Enabled = true;
 }
 
 
 Sema::CXXThisScopeRAII::~CXXThisScopeRAII() {
   if (Enabled) {
     S.CXXThisTypeOverride = OldCXXThisTypeOverride;

lib/Sema/SemaLambda.cpp

void Sema::ActOnStartOfLambdaDefinition(LambdaIntroducer &Intro,
   SmallVector<ParmVarDecl *, 8> Params;
   if (ParamInfo.getNumTypeObjects() == 0) {
     // C++11 [expr.prim.lambda]p4:
     //   If a lambda-expression does not include a lambda-declarator, it is as
     //   if the lambda-declarator were ().
     FunctionProtoType::ExtProtoInfo EPI(Context.getDefaultCallingConvention(
         /*IsVariadic=*/false, /*IsCXXMethod=*/true));
     EPI.HasTrailingReturn = true;
-    EPI.TypeQuals |= DeclSpec::TQ_const;
+    EPI.TypeQuals.addConst();
     // C++1y [expr.prim.lambda]:
     //   The lambda return type is 'auto', which is replaced by the
     //   trailing-return type if provided and/or deduced from 'return'
     //   statements
     // We don't do this before C++1y, because we don't support deduced return
     // types there.
     QualType DefaultTypeForNoTrailingReturn =
         getLangOpts().CPlusPlus14 ? Context.getAutoDeductType()
QualType Sema::getLambdaConversionFunctionResultType(
   // operator with some tweaks. The calling convention is the default free
   // function convention, and the type qualifications are lost.
   const FunctionProtoType::ExtProtoInfo CallOpExtInfo =
       CallOpProto->getExtProtoInfo();
   FunctionProtoType::ExtProtoInfo InvokerExtInfo = CallOpExtInfo;
   CallingConv CC = Context.getDefaultCallingConvention(
       CallOpProto->isVariadic(), /*IsCXXMethod=*/false);
   InvokerExtInfo.ExtInfo = InvokerExtInfo.ExtInfo.withCallingConv(CC);
-  InvokerExtInfo.TypeQuals = 0;
+  InvokerExtInfo.TypeQuals = Qualifiers();
   assert(InvokerExtInfo.RefQualifier == RQ_None &&
       "Lambda's call operator should not have a reference qualifier");
   return Context.getFunctionType(CallOpProto->getReturnType(),
                                  CallOpProto->getParamTypes(), InvokerExtInfo);
 }
 
 /// Add a lambda's conversion to function pointer, as described in
 /// C++11 [expr.prim.lambda]p6.
QualType InvokerFunctionTy = S.getLambdaConversionFunctionResultType(
       CallOperator->getType()->castAs<FunctionProtoType>());
   QualType PtrToFunctionTy = S.Context.getPointerType(InvokerFunctionTy);
 
   // Create the type of the conversion function.
   FunctionProtoType::ExtProtoInfo ConvExtInfo(
       S.Context.getDefaultCallingConvention(
       /*IsVariadic=*/false, /*IsCXXMethod=*/true));
   // The conversion function is always const.
-  ConvExtInfo.TypeQuals = Qualifiers::Const;
+  ConvExtInfo.TypeQuals = Qualifiers();
+  ConvExtInfo.TypeQuals.addConst();
   QualType ConvTy =
       S.Context.getFunctionType(PtrToFunctionTy, None, ConvExtInfo);
 
   SourceLocation Loc = IntroducerRange.getBegin();
   DeclarationName ConversionName
     = S.Context.DeclarationNames.getCXXConversionFunctionName(
         S.Context.getCanonicalType(PtrToFunctionTy));
   DeclarationNameLoc ConvNameLoc;
static void addBlockPointerConversion(Sema &S,
                                       CXXMethodDecl *CallOperator) {
   QualType FunctionTy = S.getLambdaConversionFunctionResultType(
       CallOperator->getType()->castAs<FunctionProtoType>());
   QualType BlockPtrTy = S.Context.getBlockPointerType(FunctionTy);
 
   FunctionProtoType::ExtProtoInfo ConversionEPI(
       S.Context.getDefaultCallingConvention(
           /*IsVariadic=*/false, /*IsCXXMethod=*/true));
-  ConversionEPI.TypeQuals = Qualifiers::Const;
+  ConversionEPI.TypeQuals = Qualifiers();
+  ConversionEPI.TypeQuals.addConst();
   QualType ConvTy = S.Context.getFunctionType(BlockPtrTy, None, ConversionEPI);
 
   SourceLocation Loc = IntroducerRange.getBegin();
   DeclarationName Name
     = S.Context.DeclarationNames.getCXXConversionFunctionName(
         S.Context.getCanonicalType(BlockPtrTy));
   DeclarationNameLoc NameLoc;
   NameLoc.NamedType.TInfo = S.Context.getTrivialTypeSourceInfo(BlockPtrTy, Loc);

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();
+    // FIXME: OpenCL: Need to consider address spaces
+    unsigned OldQuals = OldMethod->getTypeQualifiers().getCVRUQualifiers();
+    unsigned NewQuals = NewMethod->getTypeQualifiers().getCVRUQualifiers();
     if (!getLangOpts().CPlusPlus14 && NewMethod->isConstexpr() &&
         !isa<CXXConstructorDecl>(NewMethod))
       NewQuals |= Qualifiers::Const;
 
     // We do not allow overloading based off of '__restrict'.
     OldQuals &= ~Qualifiers::Restrict;
     NewQuals &= ~Qualifiers::Restrict;
     if (OldQuals != NewQuals)
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;
   }
 
-  unsigned FromQuals = FromFunction->getTypeQuals(),
-           ToQuals = ToFunction->getTypeQuals();
+  // FIXME: OpenCL: Need to consider address spaces
+  unsigned FromQuals = FromFunction->getTypeQuals().getCVRUQualifiers();
+  unsigned ToQuals = ToFunction->getTypeQuals().getCVRUQualifiers();
   if (FromQuals != ToQuals) {
     PDiag << ft_qualifer_mismatch << ToQuals << FromQuals;
     return;
   }
 
   // Handle exception specification differences on canonical type (in C++17
   // onwards).
   if (cast<FunctionProtoType>(FromFunction->getCanonicalTypeUnqualified())

 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();
-  QualType ImplicitParamType =  S.Context.getCVRQualifiedType(ClassType, Quals);
+  Qualifiers Quals;
+  if (isa<CXXDestructorDecl>(Method)) {
+    Quals.addConst();
+    Quals.addVolatile();
+  } else {
+    Quals = Method->getTypeQualifiers();
+  }
+
+  QualType ImplicitParamType = S.Context.getQualifiedType(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>()) {
     FromType = PT->getPointeeType();
TryObjectArgumentInitialization(Sema &S, SourceLocation Loc, QualType FromType,
 
   // Check the ref-qualifier.
   switch (Method->getRefQualifier()) {
   case RQ_None:
     // Do nothing; we don't care about lvalueness or rvalueness.
     break;
 
   case RQ_LValue:
-    if (!FromClassification.isLValue() && Quals != Qualifiers::Const) {
+    if (!FromClassification.isLValue() && !Quals.hasOnlyConst()) {
       // non-const lvalue reference cannot bind to an rvalue
       ICS.setBad(BadConversionSequence::lvalue_ref_to_rvalue, FromType,
                  ImplicitParamType);
       return ICS;
     }
     break;
 
   case RQ_RValue:
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,
QualType fnType =
       op->getRHS()->getType()->castAs<MemberPointerType>()->getPointeeType();
 
     const FunctionProtoType *proto = fnType->castAs<FunctionProtoType>();
     QualType resultType = proto->getCallResultType(Context);
     ExprValueKind valueKind = Expr::getValueKindForType(proto->getReturnType());
 
     // Check that the object type isn't more qualified than the
     // member function we're calling.
-    Qualifiers funcQuals = Qualifiers::fromCVRMask(proto->getTypeQuals());
+    Qualifiers funcQuals = proto->getTypeQuals();
 
     QualType objectType = op->getLHS()->getType();
     if (op->getOpcode() == BO_PtrMemI)
       objectType = objectType->castAs<PointerType>()->getPointeeType();
     Qualifiers objectQuals = objectType.getQualifiers();
 
     Qualifiers difference = objectQuals - funcQuals;
     difference.removeObjCGCAttr();

lib/Sema/SemaTemplate.cpp

if (FunctionTemplateDecl *FunTmpl = dyn_cast<FunctionTemplateDecl>(Ovl)) {
       // specializes), so adjust it now assuming it specializes this template.
       QualType FT = FD->getType();
       if (FD->isConstexpr()) {
         CXXMethodDecl *OldMD =
           dyn_cast<CXXMethodDecl>(FunTmpl->getTemplatedDecl());
         if (OldMD && OldMD->isConst()) {
           const FunctionProtoType *FPT = FT->castAs<FunctionProtoType>();
           FunctionProtoType::ExtProtoInfo EPI = FPT->getExtProtoInfo();
-          EPI.TypeQuals |= Qualifiers::Const;
+          EPI.TypeQuals.addConst();
           FT = Context.getFunctionType(FPT->getReturnType(),
                                        FPT->getParamTypes(), EPI);
         }
       }
 
       TemplateArgumentListInfo Args;
       if (ExplicitTemplateArgs)
         Args = *ExplicitTemplateArgs;

lib/Sema/SemaTemplateDeduction.cpp

Sema::SubstituteExplicitTemplateArguments(
   QualType ResultType;
   {
     // C++11 [expr.prim.general]p3:
     //   If a declaration declares a member function or member function
     //   template of a class X, the expression this is a prvalue of type
     //   "pointer to cv-qualifier-seq X" between the optional cv-qualifer-seq
     //   and the end of the function-definition, member-declarator, or
     //   declarator.
-    unsigned ThisTypeQuals = 0;
+    Qualifiers ThisTypeQuals;
     CXXRecordDecl *ThisContext = nullptr;
     if (CXXMethodDecl *Method = dyn_cast<CXXMethodDecl>(Function)) {
       ThisContext = Method->getParent();
       ThisTypeQuals = Method->getTypeQualifiers();
     }
 
     CXXThisScopeRAII ThisScope(*this, ThisContext, ThisTypeQuals,
                                getLangOpts().CPlusPlus11);
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, Method->getTypeQualifiers());
   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
 /// specialized as \p FT2.
 static bool isAtLeastAsSpecializedAs(Sema &S,

lib/Sema/SemaTemplateInstantiate.cpp

QualType TransformFunctionProtoType(TypeLocBuilder &TLB,
       // Call the base version; it will forward to our overridden version below.
       return inherited::TransformFunctionProtoType(TLB, TL);
     }
 
     template<typename Fn>
     QualType TransformFunctionProtoType(TypeLocBuilder &TLB,
                                         FunctionProtoTypeLoc TL,
                                         CXXRecordDecl *ThisContext,
-                                        unsigned ThisTypeQuals,
+                                        Qualifiers ThisTypeQuals,
                                         Fn TransformExceptionSpec);
 
     ParmVarDecl *TransformFunctionTypeParam(ParmVarDecl *OldParm,
                                             int indexAdjustment,
                                             Optional<unsigned> NumExpansions,
                                             bool ExpectParameterPack);
 
     /// Transforms a template type parameter type by performing
return SemaRef.BuildCXXDefaultArgExpr(E->getUsedLocation(),
                            cast<FunctionDecl>(E->getParam()->getDeclContext()),
                                         E->getParam());
 }
 
 template<typename Fn>
 QualType TemplateInstantiator::TransformFunctionProtoType(TypeLocBuilder &TLB,
                                  FunctionProtoTypeLoc TL,
                                  CXXRecordDecl *ThisContext,
-                                 unsigned ThisTypeQuals,
+                                 Qualifiers ThisTypeQuals,
                                  Fn TransformExceptionSpec) {
   // We need a local instantiation scope for this function prototype.
   LocalInstantiationScope Scope(SemaRef, /*CombineWithOuterScope=*/true);
   return inherited::TransformFunctionProtoType(
       TLB, TL, ThisContext, ThisTypeQuals, TransformExceptionSpec);
 }
 
 ParmVarDecl *

 /// type of a FunctionDecl.  Its purpose is solely to force the
 /// instantiation of default-argument expressions and to avoid
 /// instantiating an exception-specification.
 TypeSourceInfo *Sema::SubstFunctionDeclType(TypeSourceInfo *T,
                                 const MultiLevelTemplateArgumentList &Args,
                                 SourceLocation Loc,
                                 DeclarationName Entity,
                                 CXXRecordDecl *ThisContext,
-                                unsigned ThisTypeQuals) {
+                                Qualifiers ThisTypeQuals) {
   assert(!CodeSynthesisContexts.empty() &&
          "Cannot perform an instantiation without some context on the "
          "instantiation stack");
 
   if (!NeedsInstantiationAsFunctionType(T))
     return T;
 
   TemplateInstantiator Instantiator(*this, Args, Loc, Entity);
for (LateInstantiatedAttrVec::iterator I = LateAttrs.begin(),
        E = LateAttrs.end(); I != E; ++I) {
     assert(CurrentInstantiationScope == Instantiator.getStartingScope());
     CurrentInstantiationScope = I->Scope;
 
     // Allow 'this' within late-parsed attributes.
     NamedDecl *ND = dyn_cast<NamedDecl>(I->NewDecl);
     CXXRecordDecl *ThisContext =
         dyn_cast_or_null<CXXRecordDecl>(ND->getDeclContext());
-    CXXThisScopeRAII ThisScope(*this, ThisContext, /*TypeQuals*/0,
+    CXXThisScopeRAII ThisScope(*this, ThisContext, Qualifiers(),
                                ND && ND->isCXXInstanceMember());
 
     Attr *NewAttr =
       instantiateTemplateAttribute(I->TmplAttr, Context, *this, TemplateArgs);
     I->NewDecl->addAttr(NewAttr);
     LocalInstantiationScope::deleteScopes(I->Scope,
                                           Instantiator.getStartingScope());
   }
bool Sema::InstantiateInClassInitializer(
   ContextRAII SavedContext(*this, Instantiation->getParent());
   EnterExpressionEvaluationContext EvalContext(
       *this, Sema::ExpressionEvaluationContext::PotentiallyEvaluated);
 
   LocalInstantiationScope Scope(*this, true);
 
   // Instantiate the initializer.
   ActOnStartCXXInClassMemberInitializer();
-  CXXThisScopeRAII ThisScope(*this, Instantiation->getParent(), /*TypeQuals=*/0);
+  CXXThisScopeRAII ThisScope(*this, Instantiation->getParent(), Qualifiers());
 
   ExprResult NewInit = SubstInitializer(OldInit, TemplateArgs,
                                         /*CXXDirectInit=*/false);
   Expr *Init = NewInit.get();
   assert((!Init || !isa<ParenListExpr>(Init)) && "call-style init in class");
   ActOnFinishCXXInClassMemberInitializer(
       Instantiation, Init ? Init->getBeginLoc() : SourceLocation(), Init);
 

lib/Sema/SemaTemplateInstantiateDecl.cpp

if (auto *DRE = dyn_cast<DeclRefExpr>(E->IgnoreParenImpCasts()))
       if (auto *PVD = dyn_cast<ParmVarDecl>(DRE->getDecl())) {
         Sema::ContextRAII SavedContext(S, FD);
         LocalInstantiationScope Local(S);
         if (FD->getNumParams() > PVD->getFunctionScopeIndex())
           Local.InstantiatedLocal(
               PVD, FD->getParamDecl(PVD->getFunctionScopeIndex()));
         return S.SubstExpr(E, TemplateArgs);
       }
-    Sema::CXXThisScopeRAII ThisScope(S, ThisContext, /*TypeQuals=*/0,
+    Sema::CXXThisScopeRAII ThisScope(S, ThisContext, Qualifiers(),
                                      FD->isCXXInstanceMember());
     return S.SubstExpr(E, TemplateArgs);
   };
 
   ExprResult Simdlen;
   if (auto *E = Attr.getSimdlen())
     Simdlen = Subst(E);
 
void Sema::InstantiateAttrsForDecl(
     Decl *New, LateInstantiatedAttrVec *LateAttrs,
     LocalInstantiationScope *OuterMostScope) {
   if (NamedDecl *ND = dyn_cast<NamedDecl>(New)) {
     for (const auto *TmplAttr : Tmpl->attrs()) {
       // FIXME: If any of the special case versions from InstantiateAttrs become
       // applicable to template declaration, we'll need to add them here.
       CXXThisScopeRAII ThisScope(
           *this, dyn_cast_or_null<CXXRecordDecl>(ND->getDeclContext()),
-          /*TypeQuals*/ 0, ND->isCXXInstanceMember());
+          Qualifiers(), ND->isCXXInstanceMember());
 
       Attr *NewAttr = sema::instantiateTemplateAttributeForDecl(
           TmplAttr, Context, *this, TemplateArgs);
       if (NewAttr)
         New->addAttr(NewAttr);
     }
   }
 }
if (TmplAttr->isLateParsed() && LateAttrs) {
       if (CurrentInstantiationScope)
         Saved = CurrentInstantiationScope->cloneScopes(OuterMostScope);
       LateAttrs->push_back(LateInstantiatedAttribute(TmplAttr, Saved, New));
     } else {
       // Allow 'this' within late-parsed attributes.
       NamedDecl *ND = dyn_cast<NamedDecl>(New);
       CXXRecordDecl *ThisContext =
           dyn_cast_or_null<CXXRecordDecl>(ND->getDeclContext());
-      CXXThisScopeRAII ThisScope(*this, ThisContext, /*TypeQuals*/0,
+      CXXThisScopeRAII ThisScope(*this, ThisContext, Qualifiers(),
                                  ND && ND->isCXXInstanceMember());
 
       Attr *NewAttr = sema::instantiateTemplateAttribute(TmplAttr, Context,
                                                          *this, TemplateArgs);
       if (NewAttr)
         New->addAttr(NewAttr);
     }
   }
SemaRef.ActOnOpenMPDeclareReductionCombinerStart(
         /*S=*/nullptr, NewDRD);
     SemaRef.CurrentInstantiationScope->InstantiatedLocal(
         cast<DeclRefExpr>(D->getCombinerIn())->getDecl(),
         cast<DeclRefExpr>(NewDRD->getCombinerIn())->getDecl());
     SemaRef.CurrentInstantiationScope->InstantiatedLocal(
         cast<DeclRefExpr>(D->getCombinerOut())->getDecl(),
         cast<DeclRefExpr>(NewDRD->getCombinerOut())->getDecl());
     auto *ThisContext = dyn_cast_or_null<CXXRecordDecl>(Owner);
-    Sema::CXXThisScopeRAII ThisScope(SemaRef, ThisContext, /*TypeQuals*/ 0,
+    Sema::CXXThisScopeRAII ThisScope(SemaRef, ThisContext, Qualifiers(),
                                      ThisContext);
     SubstCombiner = SemaRef.SubstExpr(D->getCombiner(), TemplateArgs).get();
     SemaRef.ActOnOpenMPDeclareReductionCombinerEnd(NewDRD, SubstCombiner);
     // Initializers instantiation sequence.
     if (D->getInitializer()) {
       VarDecl *OmpPrivParm =
           SemaRef.ActOnOpenMPDeclareReductionInitializerStart(
               /*S=*/nullptr, NewDRD);

 TypeSourceInfo*
 TemplateDeclInstantiator::SubstFunctionType(FunctionDecl *D,
                               SmallVectorImpl<ParmVarDecl *> &Params) {
   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;
+  Qualifiers ThisTypeQuals;
   if (CXXMethodDecl *Method = dyn_cast<CXXMethodDecl>(D)) {
     ThisContext = cast<CXXRecordDecl>(Owner);
     ThisTypeQuals = Method->getTypeQualifiers();
   }
 
   TypeSourceInfo *NewTInfo
     = SemaRef.SubstFunctionDeclType(OldTInfo, TemplateArgs,
                                     D->getTypeSpecStartLoc(),

lib/Sema/SemaType.cpp

static QualType inferARCLifetimeForPointee(Sema &S, QualType type,
   assert(implicitLifetime && "didn't infer any lifetime!");
 
   Qualifiers qs;
   qs.addObjCLifetime(implicitLifetime);
   return S.Context.getQualifiedType(type, qs);
 }
 
 static std::string getFunctionQualifiersAsString(const FunctionProtoType *FnTy){
-  std::string Quals =
-    Qualifiers::fromCVRMask(FnTy->getTypeQuals()).getAsString();
+  std::string Quals = FnTy->getTypeQuals().getAsString();
 
   switch (FnTy->getRefQualifier()) {
   case RQ_None:
     break;
 
   case RQ_LValue:
     if (!Quals.empty())
       Quals += ' ';

 } // end anonymous namespace
 
 /// Check whether the type T is a qualified function type, and if it is,
 /// diagnose that it cannot be contained within the given kind of declarator.
 static bool checkQualifiedFunction(Sema &S, QualType T, SourceLocation Loc,
                                    QualifiedFunctionKind QFK) {
   // Does T refer to a function type with a cv-qualifier or a ref-qualifier?
   const FunctionProtoType *FPT = T->getAs<FunctionProtoType>();
-  if (!FPT || (FPT->getTypeQuals() == 0 && FPT->getRefQualifier() == RQ_None))
+  if (!FPT || (FPT->getTypeQuals().empty() && FPT->getRefQualifier() == RQ_None))
     return false;
 
   S.Diag(Loc, diag::err_compound_qualified_function_type)
     << QFK << isa<FunctionType>(T.IgnoreParens()) << T
     << getFunctionQualifiersAsString(FPT);
   return true;
 }
 

 
 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.
   DeclarationName Name;
   if (D.getIdentifier())
     Name = D.getIdentifier();
 
   // Does this declaration declare a typedef-name?
   bool IsTypedefName =
     D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_typedef ||
     D.getContext() == DeclaratorContext::AliasDeclContext ||
     D.getContext() == DeclaratorContext::AliasTemplateContext;
 
   // Does T refer to a function type with a cv-qualifier or a ref-qualifier?
   bool IsQualifiedFunction = T->isFunctionProtoType() &&
-      (T->castAs<FunctionProtoType>()->getTypeQuals() != 0 ||
+      (!T->castAs<FunctionProtoType>()->getTypeQuals().empty() ||
        T->castAs<FunctionProtoType>()->getRefQualifier() != RQ_None);
 
   // If T is 'decltype(auto)', the only declarators we can have are parens
   // and at most one function declarator if this is a function declaration.
   // If T is a deduced class template specialization type, we can have no
   // declarator chunks at all.
   if (auto *DT = T->getAs<DeducedType>()) {
     const AutoType *AT = T->getAs<AutoType>();
case DeclaratorChunk::Function: {
           T = Context.getFunctionNoProtoType(T, EI);
           break;
         }
 
         FunctionProtoType::ExtProtoInfo EPI;
         EPI.ExtInfo = EI;
         EPI.Variadic = FTI.isVariadic;
         EPI.HasTrailingReturn = FTI.hasTrailingReturnType();
-        EPI.TypeQuals = FTI.TypeQuals;
+        EPI.TypeQuals.addCVRUQualifiers(FTI.TypeQuals);
         EPI.RefQualifier = !FTI.hasRefQualifier()? RQ_None
                     : FTI.RefQualifierIsLValueRef? RQ_LValue
                     : RQ_RValue;
 
         // Otherwise, we have a function with a parameter list that is
         // potentially variadic.
         SmallVector<QualType, 16> ParamTys;
         ParamTys.reserve(FTI.NumParams);
case DeclaratorChunk::Function: {
         S.checkExceptionSpecification(D.isFunctionDeclarationContext(),
                                       FTI.getExceptionSpecType(),
                                       DynamicExceptions,
                                       DynamicExceptionRanges,
                                       NoexceptExpr,
                                       Exceptions,
                                       EPI.ExceptionSpec);
 
+        const auto &Spec = D.getCXXScopeSpec();
+        // OpenCLCPlusPlus: A class member function has an address space.
+        if (state.getSema().getLangOpts().OpenCLCPlusPlus &&
+            ((!Spec.isEmpty() &&
+              Spec.getScopeRep()->getKind() == NestedNameSpecifier::TypeSpec) ||
+             state.getDeclarator().getContext() ==
+                 DeclaratorContext::MemberContext)) {
+          LangAS CurAS = EPI.TypeQuals.getAddressSpace();
+          // If a class member function's address space is not set, set it to
+          // __generic.
+          LangAS AS =
+              (CurAS == LangAS::Default ? LangAS::opencl_generic : CurAS);
+          EPI.TypeQuals.addAddressSpace(AS);
           T = Context.getFunctionType(T, ParamTys, EPI);
+          T = state.getSema().Context.getAddrSpaceQualType(T, AS);
+        } else {
+          T = Context.getFunctionType(T, ParamTys, EPI);
+        }
       }
       break;
     }
     case DeclaratorChunk::MemberPointer: {
       // The scope spec must refer to a class, or be dependent.
       CXXScopeSpec &SS = DeclType.Mem.Scope();
       QualType ClsType;
 
if (IsQualifiedFunction &&
 
       S.Diag(Loc, diag::err_invalid_qualified_function_type)
         << Kind << D.isFunctionDeclarator() << T
         << getFunctionQualifiersAsString(FnTy)
         << FixItHint::CreateRemoval(RemovalRange);
 
       // Strip the cv-qualifiers and ref-qualifiers from the type.
       FunctionProtoType::ExtProtoInfo EPI = FnTy->getExtProtoInfo();
-      EPI.TypeQuals = 0;
+      EPI.TypeQuals.removeCVRQualifiers();
       EPI.RefQualifier = RQ_None;
 
       T = Context.getFunctionType(FnTy->getReturnType(), FnTy->getParamTypes(),
                                   EPI);
       // Rebuild any parens around the identifier in the function type.
       for (unsigned i = 0, e = D.getNumTypeObjects(); i != e; ++i) {
         if (D.getTypeObject(i).Kind != DeclaratorChunk::Paren)
           break;

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

lib/Sema/TreeTransform.h

 #define TYPELOC(CLASS, PARENT)                                   \
   QualType Transform##CLASS##Type(TypeLocBuilder &TLB, CLASS##TypeLoc T);
 #include "clang/AST/TypeLocNodes.def"
 
   template<typename Fn>
   QualType TransformFunctionProtoType(TypeLocBuilder &TLB,
                                       FunctionProtoTypeLoc TL,
                                       CXXRecordDecl *ThisContext,
-                                      unsigned ThisTypeQuals,
+                                      Qualifiers ThisTypeQuals,
                                       Fn TransformExceptionSpec);
 
   bool TransformExceptionSpec(SourceLocation Loc,
                               FunctionProtoType::ExceptionSpecInfo &ESI,
                               SmallVectorImpl<QualType> &Exceptions,
                               bool &Changed);
 
   StmtResult TransformSEHHandler(Stmt *Handler);
if (((T.getAddressSpace() != LangAS::Default &&
     SemaRef.Diag(Loc, diag::err_address_space_mismatch_templ_inst)
         << TL.getType() << T;
     return QualType();
   }
 
   // C++ [dcl.fct]p7:
   //   [When] adding cv-qualifications on top of the function type [...] the
   //   cv-qualifiers are ignored.
-  if (T->isFunctionType())
+  if (T->isFunctionType()) {
+    T = SemaRef.getASTContext().getAddrSpaceQualType(T,
+                                                     Quals.getAddressSpace());
     return T;
+  }
 
   // C++ [dcl.ref]p1:
   //   when the cv-qualifiers are introduced through the use of a typedef-name
   //   or decltype-specifier [...] the cv-qualifiers are ignored.
   // Note that [dcl.ref]p1 lists all cases in which cv-qualifiers can be
   // applied to a reference type.
   if (T->isReferenceType()) {
     // The only qualifier that applies to a reference type is restrict.

 
 template<typename Derived>
 QualType
 TreeTransform<Derived>::TransformFunctionProtoType(TypeLocBuilder &TLB,
                                                    FunctionProtoTypeLoc TL) {
   SmallVector<QualType, 4> ExceptionStorage;
   TreeTransform *This = this; // Work around gcc.gnu.org/PR56135.
   return getDerived().TransformFunctionProtoType(
-      TLB, TL, nullptr, 0,
+      TLB, TL, nullptr, Qualifiers(),
       [&](FunctionProtoType::ExceptionSpecInfo &ESI, bool &Changed) {
         return This->TransformExceptionSpec(TL.getBeginLoc(), ESI,
                                             ExceptionStorage, Changed);
       });
 }
 
 template<typename Derived> template<typename Fn>
 QualType TreeTransform<Derived>::TransformFunctionProtoType(
     TypeLocBuilder &TLB, FunctionProtoTypeLoc TL, CXXRecordDecl *ThisContext,
-    unsigned ThisTypeQuals, Fn TransformExceptionSpec) {
+    Qualifiers ThisTypeQuals, Fn TransformExceptionSpec) {
 
   // Transform the parameters and return type.
   //
   // We are required to instantiate the params and return type in source order.
   // When the function has a trailing return type, we instantiate the
   // parameters before the return type,  since the return type can then refer
   // to the parameters themselves (via decltype, sizeof, etc.).
   //
   SmallVector<QualType, 4> ParamTypes;
   SmallVector<ParmVarDecl*, 4> ParamDecls;
   Sema::ExtParameterInfoBuilder ExtParamInfos;
   const FunctionProtoType *T = TL.getTypePtr();
 
   QualType ResultType;
 
   if (T->hasTrailingReturn()) {
     if (getDerived().TransformFunctionTypeParams(
             TL.getBeginLoc(), TL.getParams(),
             TL.getTypePtr()->param_type_begin(),
             T->getExtParameterInfosOrNull(),
             ParamTypes, &ParamDecls, ExtParamInfos))
       return QualType();
 
     {
       // C++11 [expr.prim.general]p3:
       //   If a declaration declares a member function or member function
TypeSourceInfo *NewCallOpTSI = nullptr;
     TypeSourceInfo *OldCallOpTSI = E->getCallOperator()->getTypeSourceInfo();
     FunctionProtoTypeLoc OldCallOpFPTL =
         OldCallOpTSI->getTypeLoc().getAs<FunctionProtoTypeLoc>();
 
     TypeLocBuilder NewCallOpTLBuilder;
     SmallVector<QualType, 4> ExceptionStorage;
     TreeTransform *This = this; // Work around gcc.gnu.org/PR56135.
     QualType NewCallOpType = TransformFunctionProtoType(
-        NewCallOpTLBuilder, OldCallOpFPTL, nullptr, 0,
+        NewCallOpTLBuilder, OldCallOpFPTL, nullptr, Qualifiers(),
         [&](FunctionProtoType::ExceptionSpecInfo &ESI, bool &Changed) {
           return This->TransformExceptionSpec(OldCallOpFPTL.getBeginLoc(), ESI,
                                               ExceptionStorage, Changed);
         });
     if (NewCallOpType.isNull())
       return ExprError();
     NewCallOpTSI = NewCallOpTLBuilder.getTypeSourceInfo(getSema().Context,
                                                         NewCallOpType);

lib/Serialization/ASTReader.cpp

EPI.ExtInfo = FunctionType::ExtInfo(/*noreturn*/ Record[1],
                                         /*produces*/ Record[5],
                                         /*nocallersavedregs*/ Record[6],
                                         /*nocfcheck*/ Record[7]);
 
     unsigned Idx = 8;
 
     EPI.Variadic = Record[Idx++];
     EPI.HasTrailingReturn = Record[Idx++];
-    EPI.TypeQuals = Record[Idx++];
+    EPI.TypeQuals = Qualifiers::fromOpaqueValue(Record[Idx++]);
     EPI.RefQualifier = static_cast<RefQualifierKind>(Record[Idx++]);
     SmallVector<QualType, 8> ExceptionStorage;
     readExceptionSpec(*Loc.F, ExceptionStorage, EPI.ExceptionSpec, Record, Idx);
 
     unsigned NumParams = Record[Idx++];
     SmallVector<QualType, 16> ParamTypes;
     for (unsigned I = 0; I != NumParams; ++I)
       ParamTypes.push_back(readType(*Loc.F, Record, Idx));

lib/Serialization/ASTWriter.cpp

static void addExceptionSpec(const FunctionProtoType *T,
   }
 }
 
 void ASTTypeWriter::VisitFunctionProtoType(const FunctionProtoType *T) {
   VisitFunctionType(T);
 
   Record.push_back(T->isVariadic());
   Record.push_back(T->hasTrailingReturn());
-  Record.push_back(T->getTypeQuals());
+  Record.push_back(T->getTypeQuals().getAsOpaqueValue());
   Record.push_back(static_cast<unsigned>(T->getRefQualifier()));
   addExceptionSpec(T, Record);
 
   Record.push_back(T->getNumParams());
   for (unsigned I = 0, N = T->getNumParams(); I != N; ++I)
     Record.AddTypeRef(T->getParamType(I));
 
   if (T->hasExtParameterInfos()) {

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; }
+  // FIXME: Does not work yet.
+  // C(C &&c) { v = c.v; }
+  C(const C &c) { v = c.v; }
+  C &operator=(const C &c) {
+    v = c.v;
+    return *this;
+  }
+  // FIXME: Does not work yet.
+  //C &operator=(C&& c) & {
+  //  v = c.v;
+  //  return *this;
+  //}
+
+  int get() { return v; }
+
+  int outside();
+};
+
+int C::outside() {
+  return v;
+}
+
+extern C&& foo();
+
+__global C c;
+
+__kernel void test__global() {
+  int i = c.get();
+  int i2 = c.outside();
+  C c1(c);
+  C c2;
+  c2 = c1;
+  // FIXME: Does not work yet.
+  // C c3 = c1 + c2;
+  // C c4(foo());
+  // C c5 = foo();
+
+}
+
+// CHECK-LABEL: @__cxx_global_var_init()
+// CHECK: call void @_ZNU3AS41CC1Ev(%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: @_ZNU3AS41CC1Ev(%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 @_ZNU3AS41CC2Ev(%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 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+
+// Test the address space of 'this' when invoking a method that is declared in the file contex.
+// CHECK: %call1 = call i32 @_ZNU3AS41C7outsideEv(%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 @_ZNU3AS41CC1ERU3AS4KS_(%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 @_ZNU3AS41CC1Ev(%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:   %call2 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%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 @_ZNU3AS41CC1Ev(%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 @_ZNU3AS41C3getEv(%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 @_ZNU3AS41CC1Ev(%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)* @_ZNU3AS41CaSERU3AS4KS_(%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 @_ZNU3AS41C3getEv(%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 @_ZNU3AS41CC1ERU3AS4KS_(%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 @_ZNU3AS41CC1Ev(%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)* @_ZNU3AS41CaSERU3AS4KS_(%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 @_ZNU3AS41C3getEv(%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 @_ZNU3AS41CC1ERU3AS4KS_(%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 @_ZNU3AS41CC1Ev(%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)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %5)

test/CodeGenOpenCLCXX/template-address-spaces.cl

 // RUN: %clang_cc1 -cl-std=c++ %s -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck %s
 
 template <typename T>
 struct S{
   T a;
   T foo();
 };
 
 template<typename T>
 T S<T>::foo() { return a;}
 
-//CHECK: %struct.S = type { i32 }
-//CHECK: %struct.S.0 = type { i32 addrspace(4)* }
-//CHECK: %struct.S.1 = type { i32 addrspace(1)* }
+// CHECK: %struct.S = type { i32 }
+// CHECK: %struct.S.0 = type { i32 addrspace(4)* }
+// CHECK: %struct.S.1 = type { i32 addrspace(1)* }
 
-//CHECK: i32 @_ZN1SIiE3fooEv(%struct.S* %this)
-//CHECK: i32 addrspace(4)* @_ZN1SIPU3AS4iE3fooEv(%struct.S.0* %this)
-//CHECK: i32 addrspace(1)* @_ZN1SIPU3AS1iE3fooEv(%struct.S.1* %this)
+// CHECK:  %0 = addrspacecast %struct.S* %sint to %struct.S addrspace(4)*
+// CHECK:  %call = call i32 @_ZNU3AS41SIiE3fooEv(%struct.S addrspace(4)* %0) #1
+// CHECK:  %1 = addrspacecast %struct.S.0* %sintptr to %struct.S.0 addrspace(4)*
+// CHECK:  %call1 = call i32 addrspace(4)* @_ZNU3AS41SIPU3AS4iE3fooEv(%struct.S.0 addrspace(4)* %1) #1
+// CHECK:  %2 = addrspacecast %struct.S.1* %sintptrgl to %struct.S.1 addrspace(4)*
+// CHECK:  %call2 = call i32 addrspace(1)* @_ZNU3AS41SIPU3AS1iE3fooEv(%struct.S.1 addrspace(4)* %2) #1
 
 void bar(){
   S<int> sint;
   S<int*> sintptr;
   S<__global int*> sintptrgl;
 
   sint.foo();
   sintptr.foo();
   sintptrgl.foo();
 }

test/SemaOpenCLCXX/address-space-templates.cl

 //RUN: %clang_cc1 %s -cl-std=c++ -pedantic -verify -fsyntax-only
 
 template <typename T>
 struct S {
   T a;        // expected-error{{field may not be qualified with an address space}}
   T f1();     // expected-error{{function type may not be qualified with an address space}}
-  void f2(T); // expected-error{{parameter may not be qualified with an address space}}
+  // FIXME: Should only get the error message once.
+  void f2(T); // expected-error{{parameter may not be qualified with an address space}} expected-error{{parameter may not be qualified with an address space}}
+
 };
 
 template <typename T>
 T foo1(__global T *i) { // expected-note{{candidate template ignored: substitution failure [with T = __local int]: conflicting address space qualifiers are provided between types '__global T' and '__local int'}}
   return *i;
 }
 
 template <typename T>

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 =