[OpenCL] Improve printing and semantic check related to implicit addr space

[yaxunl]

Summary

There are two issues:

1. only (void*)0 should be treated as nullptr

2. only explicit addr space should be printed

This patch introduces a flag in Qualifier to indicating a non-default address space qualifier is deduced by context. Only
non-implicit address space qualifier will be print out when printing AST. It is also used to identify nullptr.

However this review does not rule out alternative approaches, e.g. using AttributedType. We will explore alternative approaches.

This patch depends on https://reviews.llvm.org/D35082

[Anastasia]

LGTM! Thanks!

Can we close https://bugs.llvm.org/show_bug.cgi?id=33418 after this commit?

test/SemaOpenCL/null_literal.cl

38

Does this extra cast test anything we already miss to test?

[yaxunl]

In D38857#896994, @Anastasia wrote:

LGTM! Thanks!

Can we close https://bugs.llvm.org/show_bug.cgi?id=33418 after this commit?

Will do.

test/SemaOpenCL/null_literal.cl

38

Yes. It tests a generic pointer of zero value can be explicitly casted to a global pointer. This should be true for any generic pointer, however since pointers with zero value have special handling, we want to make sure this still works.

[rjmccall]

include/clang/AST/Type.h

337

This is probably cleaner as:

Mask = (Value ? (Mask | ImplicitAddrSpaceMask) : (Mask & ~ImplicitAddrSpaceMask));

lib/AST/ASTContext.cpp

2290

It looks like your changes here are making implicitness part of the canonical type, which is wrong, because implicitly- and explicitly-qualified types are not actually different types.

That is fixable, but I'm going to ask you to investigate whether you can solve this problem with AttributedType before you introduce this complexity into the qualifier system.

include/clang/AST/ASTContext.h

 
 public:
   /// \brief Return the uniqued reference to the type for an address space
   /// qualified type with the specified type and address space.
   ///
   /// The resulting type has a union of the qualifiers from T and the address
   /// space. If T already has an address space specifier, it is silently
   /// replaced.
-  QualType getAddrSpaceQualType(QualType T, unsigned AddressSpace) const;
+  QualType getAddrSpaceQualType(QualType T, unsigned AddressSpace,
+                                bool IsImplicit = false) const;
 
   /// \brief Remove any existing address space on the type and returns the type
   /// with qualifiers intact (or that's the idea anyway)
   ///
   /// The return type should be T with all prior qualifiers minus the address
   /// space.
   QualType removeAddrSpaceQualType(QualType T) const;
 

include/clang/AST/Type.h

enum ObjCLifetime {
     OCL_Weak,
 
     /// Assigning into this object requires a lifetime extension.
     OCL_Autoreleasing
   };
 
   enum {
     /// The maximum supported address space number.
-    /// 23 bits should be enough for anyone.
-    MaxAddressSpace = 0x7fffffu,
+    /// 22 bits should be enough for anyone.
+    MaxAddressSpace = 0x3fffffu,
 
     /// The width of the "fast" qualifier mask.
     FastWidth = 3,
 
     /// The fast qualifier mask.
     FastMask = (1 << FastWidth) - 1
   };
 
public:
   }
 
   /// True if the lifetime is either strong or weak.
   bool hasStrongOrWeakObjCLifetime() const {
     ObjCLifetime lifetime = getObjCLifetime();
     return (lifetime == OCL_Strong || lifetime == OCL_Weak);
   }
 
+  /// True if the non-default address space is not explicit in the source
+  /// code but deduced by context. This flag is used when printing
+  /// types or performing semantic checks if the explicity of an address
+  /// space makes difference.
+  bool isAddressSpaceImplicit() const { return Mask & ImplicitAddrSpaceMask; }
+  void setAddressSpaceImplicit(bool Value) {
+    Mask = (Mask & ~ImplicitAddrSpaceMask) |
+           (((uint32_t)Value) << ImplicitAddrSpaceShift);
+  }
+  void removeImplicitAddressSpaceFlag() { setAddressSpaceImplicit(false); }
   bool hasAddressSpace() const { return Mask & AddressSpaceMask; }
   unsigned getAddressSpace() const { return Mask >> AddressSpaceShift; }
   bool hasTargetSpecificAddressSpace() const {
     return getAddressSpace() >= LangAS::FirstTargetAddressSpace;
   }
   /// Get the address space attribute value to be printed by diagnostics.
   unsigned getAddressSpaceAttributePrintValue() const {
     auto Addr = getAddressSpace();
     // This function is not supposed to be used with language specific
     // address spaces. If that happens, the diagnostic message should consider
     // printing the QualType instead of the address space value.
     assert(Addr == 0 || hasTargetSpecificAddressSpace());
     if (Addr)
       return Addr - LangAS::FirstTargetAddressSpace;
     // TODO: The diagnostic messages where Addr may be 0 should be fixed
     // since it cannot differentiate the situation where 0 denotes the default
     // address space or user specified __attribute__((address_space(0))).
     return 0;
   }
   void setAddressSpace(unsigned space) {
     assert(space <= MaxAddressSpace);
     Mask = (Mask & ~AddressSpaceMask)
          | (((uint32_t) space) << AddressSpaceShift);
   }
-  void removeAddressSpace() { setAddressSpace(0); }
+  void removeAddressSpace() {
+    setAddressSpace(0);
+    removeImplicitAddressSpaceFlag();
+  }
   void addAddressSpace(unsigned space) {
     assert(space);
     setAddressSpace(space);
   }
 
   // Fast qualifiers are those that can be allocated directly
   // on a QualType object.
   bool hasFastQualifiers() const { return getFastQualifiers(); }
public:
   void print(raw_ostream &OS, const PrintingPolicy &Policy,
              bool appendSpaceIfNonEmpty = false) const;
 
   void Profile(llvm::FoldingSetNodeID &ID) const {
     ID.AddInteger(Mask);
   }
 
 private:
-
-  // bits:     |0 1 2|3|4 .. 5|6  ..  8|9   ...   31|
-  //           |C R V|U|GCAttr|Lifetime|AddressSpace|
+  // bits:     |0 1 2|3|4 .. 5|6  ..  8|9|10   ...   31|
+  //           |C R V|U|GCAttr|Lifetime|I|AddressSpace |
+  // I is abbreviation for ImplicitAddrSpace.
   uint32_t Mask = 0;
 
   static const uint32_t UMask = 0x8;
   static const uint32_t UShift = 3;
   static const uint32_t GCAttrMask = 0x30;
   static const uint32_t GCAttrShift = 4;
   static const uint32_t LifetimeMask = 0x1C0;
   static const uint32_t LifetimeShift = 6;
+  static const uint32_t ImplicitAddrSpaceMask = 0x200;
+  static const uint32_t ImplicitAddrSpaceShift = 9;
   static const uint32_t AddressSpaceMask =
-      ~(CVRMask | UMask | GCAttrMask | LifetimeMask);
-  static const uint32_t AddressSpaceShift = 9;
+      ~(CVRMask | UMask | GCAttrMask | LifetimeMask | ImplicitAddrSpaceMask);
+  static const uint32_t AddressSpaceShift = 10;
 };
 
 /// A std::pair-like structure for storing a qualified type split
 /// into its local qualifiers and its locally-unqualified type.
 struct SplitQualType {
   /// The locally-unqualified type.
   const Type *Ty;
 

lib/AST/ASTContext.cpp

     (void) ExtQualNodes.FindNodeOrInsertPos(ID, insertPos);
   }
 
   ExtQuals *eq = new (*this, TypeAlignment) ExtQuals(baseType, canon, quals);
   ExtQualNodes.InsertNode(eq, insertPos);
   return QualType(eq, fastQuals);
 }
 
-QualType
-ASTContext::getAddrSpaceQualType(QualType T, unsigned AddressSpace) const {
+QualType ASTContext::getAddrSpaceQualType(QualType T, unsigned AddressSpace,
+                                          bool ImplicitFlag) const {
   QualType CanT = getCanonicalType(T);
-  if (CanT.getAddressSpace() == AddressSpace)
+  if (CanT.getAddressSpace() == AddressSpace &&
+      CanT.getQualifiers().isAddressSpaceImplicit() == ImplicitFlag)
     return T;
 
   // If we are composing extended qualifiers together, merge together
   // into one ExtQuals node.
   QualifierCollector Quals;
   const Type *TypeNode = Quals.strip(T);
 
   // If this type already has an address space specified, it cannot get
   // another one.
   assert(!Quals.hasAddressSpace() &&
          "Type cannot be in multiple addr spaces!");
   Quals.addAddressSpace(AddressSpace);
+  Quals.setAddressSpaceImplicit(ImplicitFlag);
 
   return getExtQualType(TypeNode, Quals);
 }
 
 QualType ASTContext::removeAddrSpaceQualType(QualType T) const {
   // If we are composing extended qualifiers together, merge together
   // into one ExtQuals node.
   QualifierCollector Quals;

 
   // If two types are identical, they are compatible.
   if (LHSCan == RHSCan)
     return LHS;
 
   // If the qualifiers are different, the types aren't compatible... mostly.
   Qualifiers LQuals = LHSCan.getLocalQualifiers();
   Qualifiers RQuals = RHSCan.getLocalQualifiers();
+  RQuals.setAddressSpaceImplicit(LQuals.isAddressSpaceImplicit());
   if (LQuals != RQuals) {
     // If any of these qualifiers are different, we have a type
     // mismatch.
     if (LQuals.getCVRQualifiers() != RQuals.getCVRQualifiers() ||
         LQuals.getAddressSpace() != RQuals.getAddressSpace() ||
         LQuals.getObjCLifetime() != RQuals.getObjCLifetime() ||
         LQuals.hasUnaligned() != RQuals.hasUnaligned())
       return QualType();

lib/AST/Expr.cpp

       if (const PointerType *PT = CE->getType()->getAs<PointerType>()) {
         QualType Pointee = PT->getPointeeType();
         // Only (void*)0 or equivalent are treated as nullptr. If pointee type
         // has non-default address space it is not treated as nullptr.
         // (__generic void*)0 in OpenCL 2.0 should not be treated as nullptr
         // since it cannot be assigned to a pointer to constant address space.
         bool PointeeHasDefaultAS =
             Pointee.getAddressSpace() == LangAS::Default ||
-            (Ctx.getLangOpts().OpenCLVersion >= 200 &&
-             Pointee.getAddressSpace() == LangAS::opencl_generic) ||
-            (Ctx.getLangOpts().OpenCL &&
-             Ctx.getLangOpts().OpenCLVersion < 200 &&
-             Pointee.getAddressSpace() == LangAS::opencl_private);
+            Pointee.getQualifiers().isAddressSpaceImplicit();
 
         if (PointeeHasDefaultAS && Pointee->isVoidType() && // to void*
             CE->getSubExpr()->getType()->isIntegerType())   // from int.
           return CE->getSubExpr()->isNullPointerConstant(Ctx, NPC);
       }
     }
   } else if (const ImplicitCastExpr *ICE = dyn_cast<ImplicitCastExpr>(this)) {
     // Ignore the ImplicitCastExpr type entirely.

lib/AST/TypePrinter.cpp

     addSpace = true;
   }
   if (hasUnaligned()) {
     if (addSpace)
       OS << ' ';
     OS << "__unaligned";
     addSpace = true;
   }
-  if (unsigned addrspace = getAddressSpace()) {
-    if (addrspace != LangAS::opencl_private) {
+  if (!isAddressSpaceImplicit()) {
+    if (unsigned addrspace = getAddressSpace()) {
       if (addSpace)
         OS << ' ';
       addSpace = true;
       switch (addrspace) {
       case LangAS::opencl_global:
         OS << "__global";
         break;
       case LangAS::opencl_local:
         OS << "__local";
         break;
       case LangAS::opencl_constant:
       case LangAS::cuda_constant:
         OS << "__constant";
         break;
+      case LangAS::opencl_private:
+        OS << "__private";
+        break;
       case LangAS::opencl_generic:
         OS << "__generic";
         break;
       case LangAS::cuda_device:
         OS << "__device";
         break;
       case LangAS::cuda_shared:
         OS << "__shared";

lib/Sema/SemaType.cpp

             D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_extern) {
           ImpAddr = LangAS::opencl_global;
         } else {
           ImpAddr = LangAS::opencl_private;
         }
       }
     }
   }
-  T = State.getSema().Context.getAddrSpaceQualType(T, ImpAddr);
+  T = State.getSema().Context.getAddrSpaceQualType(T, ImpAddr, true);
 }
 
 static void processTypeAttrs(TypeProcessingState &state, QualType &type,
                              TypeAttrLocation TAL, AttributeList *attrs) {
   // Scan through and apply attributes to this type where it makes sense.  Some
   // attributes (such as __address_space__, __vector_size__, etc) apply to the
   // type, but others can be present in the type specifiers even though they
   // apply to the decl.  Here we apply type attributes and ignore the rest.

test/SemaOpenCL/address-spaces-conversions-cl2.0.cl

#endif
 
   AS int *var_init3 = arg_const;
 #ifndef CONSTANT
 // expected-error-re@-2{{initializing '__{{global|generic}} int *' with an expression of type '__constant int *' changes address space of pointer}}
 #endif
 
   AS int *var_init4 = arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{initializing '__{{global|constant}} int *' with an expression of type 'int *' changes address space of pointer}}
+// expected-error-re@-2{{initializing '__{{global|constant}} int *' with an expression of type '__private int *' changes address space of pointer}}
 #endif
 
   AS int *var_init5 = arg_gen;
 #ifndef GENERIC
 // expected-error-re@-2{{initializing '__{{global|constant}} int *' with an expression of type '__generic int *' changes address space of pointer}}
 #endif
 
   AS int *var_cast1 = (AS int *)arg_glob;
 #ifdef CONSTANT
 // expected-error@-2{{casting '__global int *' to type '__constant int *' changes address space of pointer}}
 #endif
 
   AS int *var_cast2 = (AS int *)arg_loc;
 #ifndef GENERIC
 // expected-error-re@-2{{casting '__local int *' to type '__{{global|constant}} int *' changes address space of pointer}}
 #endif
 
   AS int *var_cast3 = (AS int *)arg_const;
 #ifndef CONSTANT
 // expected-error-re@-2{{casting '__constant int *' to type '__{{global|generic}} int *' changes address space of pointer}}
 #endif
 
   AS int *var_cast4 = (AS int *)arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-2{{casting '__private int *' to type '__{{global|constant}} int *' changes address space of pointer}}
 #endif
 
   AS int *var_cast5 = (AS int *)arg_gen;
 #ifdef CONSTANT
 // expected-error@-2{{casting '__generic int *' to type '__constant int *' changes address space of pointer}}
 #endif
 
   AS int *var_impl;
#endif
 
   var_impl = arg_const;
 #ifndef CONSTANT
 // expected-error-re@-2{{assigning '__constant int *' to '__{{global|generic}} int *' changes address space of pointer}}
 #endif
 
   var_impl = arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{assigning 'int *' to '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-2{{assigning '__private int *' to '__{{global|constant}} int *' changes address space of pointer}}
 #endif
 
   var_impl = arg_gen;
 #ifndef GENERIC
 // expected-error-re@-2{{assigning '__generic int *' to '__{{global|constant}} int *' changes address space of pointer}}
 #endif
 
   var_cast1 = (AS int *)arg_glob;
 #ifdef CONSTANT
 // expected-error@-2{{casting '__global int *' to type '__constant int *' changes address space of pointer}}
 #endif
 
   var_cast2 = (AS int *)arg_loc;
 #ifndef GENERIC
 // expected-error-re@-2{{casting '__local int *' to type '__{{global|constant}} int *' changes address space of pointer}}
 #endif
 
   var_cast3 = (AS int *)arg_const;
 #ifndef CONSTANT
 // expected-error-re@-2{{casting '__constant int *' to type '__{{global|generic}} int *' changes address space of pointer}}
 #endif
 
   var_cast4 = (AS int *)arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-2{{casting '__private int *' to type '__{{global|constant}} int *' changes address space of pointer}}
 #endif
 
   var_cast5 = (AS int *)arg_gen;
 #ifdef CONSTANT
 // expected-error@-2{{casting '__generic int *' to type '__constant int *' changes address space of pointer}}
 #endif
 
   AS int *var_cmp;
#endif
 
   b = var_cmp == arg_const;
 #ifndef CONSTANT
 // expected-error-re@-2{{comparison between  ('__{{global|generic}} int *' and '__constant int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   b = var_cmp <= arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{comparison between  ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-2{{comparison between  ('__{{global|constant}} int *' and '__private int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   b = var_cmp >= arg_gen;
 #ifdef CONSTANT
 // expected-error@-2{{comparison between  ('__constant int *' and '__generic int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   AS int *var_sub;
#endif
 
   b = var_sub - arg_const;
 #ifndef CONSTANT
 // expected-error-re@-2{{arithmetic operation with operands of type  ('__{{global|generic}} int *' and '__constant int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   b = var_sub - arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{arithmetic operation with operands of type  ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-2{{arithmetic operation with operands of type  ('__{{global|constant}} int *' and '__private int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   b = var_sub - arg_gen;
 #ifdef CONSTANT
 // expected-error@-2{{arithmetic operation with operands of type  ('__constant int *' and '__generic int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   f_glob(var_sub);
 #ifndef GLOBAL
 // expected-error-re@-2{{passing '__{{constant|generic}} int *' to parameter of type '__global int *' changes address space of pointer}}
 #endif
 
   f_loc(var_sub); // expected-error-re{{passing '__{{global|constant|generic}} int *' to parameter of type '__local int *' changes address space of pointer}}
 
   f_const(var_sub);
 #ifndef CONSTANT
 // expected-error-re@-2{{passing '__{{global|generic}} int *' to parameter of type '__constant int *' changes address space of pointer}}
 #endif
 
-  f_priv(var_sub); // expected-error-re{{passing '__{{global|constant|generic}} int *' to parameter of type 'int *' changes address space of pointer}}
+  f_priv(var_sub); // expected-error-re{{passing '__{{global|constant|generic}} int *' to parameter of type '__private int *' changes address space of pointer}}
 
   f_gen(var_sub);
 #ifdef CONSTANT
 // expected-error@-2{{passing '__constant int *' to parameter of type '__generic int *' changes address space of pointer}}
 #endif
 }
 
 void test_ternary() {
#endif
   var_cond = 0 ? var_cond : var_const;
 #ifndef CONSTANT
 // expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|generic}} int *' and '__constant int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   private int *var_priv;
   var_gen = 0 ? var_cond : var_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and '__private int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   var_gen = 0 ? var_cond : var_gen;
 #ifdef CONSTANT
 // expected-error@-2{{conditional operator with the second and third operands of type  ('__constant int *' and '__generic int *') which are pointers to non-overlapping address spaces}}
 #endif
 
   void *var_void_gen;

 // expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|generic}} int *' and '__constant char *') which are pointers to non-overlapping address spaces}}
 #else
 // expected-warning@-4{{pointer type mismatch ('__constant int *' and '__constant char *')}}
 #endif
 
   private char *var_priv_ch;
   var_void_gen = 0 ? var_cond : var_priv_ch;
 #ifndef GENERIC
-// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and 'char *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and '__private char *') which are pointers to non-overlapping address spaces}}
 #else
-// expected-warning@-4{{pointer type mismatch ('__generic int *' and 'char *')}}
+// expected-warning@-4{{pointer type mismatch ('__generic int *' and '__private char *')}}
 #endif
 
   generic char *var_gen_ch;
   var_void_gen = 0 ? var_cond : var_gen_ch;
 #ifdef CONSTANT
 // expected-error@-2{{conditional operator with the second and third operands of type  ('__constant int *' and '__generic char *') which are pointers to non-overlapping address spaces}}
 #else
 // expected-warning-re@-4{{pointer type mismatch ('__{{global|generic}} int *' and '__generic char *')}}

test/SemaOpenCL/address-spaces.cl

__kernel void foo(__global int *gip) {
   int *ip;
 #if __OPENCL_C_VERSION__ < 200
   ip = gip; // expected-error {{assigning '__global int *' to 'int *' changes address space of pointer}}
   ip = &li; // expected-error {{assigning '__local int *' to 'int *' changes address space of pointer}}
   ip = &ci; // expected-error {{assigning '__constant int *' to 'int *' changes address space of pointer}}
 #else
   ip = gip;
   ip = &li;
-  ip = &ci; // expected-error {{assigning '__constant int *' to '__generic int *' changes address space of pointer}}
+  ip = &ci; // expected-error {{assigning '__constant int *' to 'int *' changes address space of pointer}}
 #endif
 }
 
 void explicit_cast(global int* g, local int* l, constant int* c, private int* p, const constant int *cc)
 {
   g = (global int*) l;    // expected-error {{casting '__local int *' to type '__global int *' changes address space of pointer}}
   g = (global int*) c;    // expected-error {{casting '__constant int *' to type '__global int *' changes address space of pointer}}
   g = (global int*) cc;   // expected-error {{casting 'const __constant int *' to type '__global int *' changes address space of pointer}}
-  g = (global int*) p;    // expected-error {{casting 'int *' to type '__global int *' changes address space of pointer}}
+  g = (global int*) p;    // expected-error {{casting '__private int *' to type '__global int *' changes address space of pointer}}
 
   l = (local int*) g;     // expected-error {{casting '__global int *' to type '__local int *' changes address space of pointer}}
   l = (local int*) c;     // expected-error {{casting '__constant int *' to type '__local int *' changes address space of pointer}}
   l = (local int*) cc;    // expected-error {{casting 'const __constant int *' to type '__local int *' changes address space of pointer}}
-  l = (local int*) p;     // expected-error {{casting 'int *' to type '__local int *' changes address space of pointer}}
+  l = (local int*) p;     // expected-error {{casting '__private int *' to type '__local int *' changes address space of pointer}}
 
   c = (constant int*) g;  // expected-error {{casting '__global int *' to type '__constant int *' changes address space of pointer}}
   c = (constant int*) l;  // expected-error {{casting '__local int *' to type '__constant int *' changes address space of pointer}}
-  c = (constant int*) p;  // expected-error {{casting 'int *' to type '__constant int *' changes address space of pointer}}
+  c = (constant int*) p;  // expected-error {{casting '__private int *' to type '__constant int *' changes address space of pointer}}
 
-  p = (private int*) g;   // expected-error {{casting '__global int *' to type 'int *' changes address space of pointer}}
-  p = (private int*) l;   // expected-error {{casting '__local int *' to type 'int *' changes address space of pointer}}
-  p = (private int*) c;   // expected-error {{casting '__constant int *' to type 'int *' changes address space of pointer}}
-  p = (private int*) cc;  // expected-error {{casting 'const __constant int *' to type 'int *' changes address space of pointer}}
+  p = (private int*) g;   // expected-error {{casting '__global int *' to type '__private int *' changes address space of pointer}}
+  p = (private int*) l;   // expected-error {{casting '__local int *' to type '__private int *' changes address space of pointer}}
+  p = (private int*) c;   // expected-error {{casting '__constant int *' to type '__private int *' changes address space of pointer}}
+  p = (private int*) cc;  // expected-error {{casting 'const __constant int *' to type '__private int *' changes address space of pointer}}
 }
 
 void ok_explicit_casts(global int *g, global int* g2, local int* l, local int* l2, private int* p, private int* p2)
 {
   g = (global int*) g2;
   l = (local int*) l2;
   p = (private int*) p2;
 }

test/SemaOpenCL/atomic-ops.cl

 
 atomic_int gn;
 void f(atomic_int *i, const atomic_int *ci,
        atomic_intptr_t *p, atomic_float *d,
        int *I, const int *CI,
        intptr_t *P, float *D, struct S *s1, struct S *s2,
        global atomic_int *i_g, local atomic_int *i_l, private atomic_int *i_p,
        constant atomic_int *i_c) {
-  __opencl_atomic_init(I, 5); // expected-error {{address argument to atomic operation must be a pointer to _Atomic type ('__generic int *' invalid)}}
-  __opencl_atomic_init(ci, 5); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+  __opencl_atomic_init(I, 5); // expected-error {{address argument to atomic operation must be a pointer to _Atomic type ('int *' invalid)}}
+  __opencl_atomic_init(ci, 5); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const atomic_int *' (aka 'const _Atomic(int) *') invalid)}}
 
   __opencl_atomic_load(0); // expected-error {{too few arguments to function call, expected 3, have 1}}
   __opencl_atomic_load(0, 0, 0, 0); // expected-error {{too many arguments to function call, expected 3, have 4}}
   __opencl_atomic_store(0,0,0,0); // expected-error {{address argument to atomic builtin must be a pointer}}
-  __opencl_atomic_store((int *)0, 0, 0, 0); // expected-error {{address argument to atomic operation must be a pointer to _Atomic type ('__generic int *' invalid)}}
+  __opencl_atomic_store((int *)0, 0, 0, 0); // expected-error {{address argument to atomic operation must be a pointer to _Atomic type ('int *' invalid)}}
   __opencl_atomic_store(i, 0, memory_order_relaxed, memory_scope_work_group);
-  __opencl_atomic_store(ci, 0, memory_order_relaxed, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+  __opencl_atomic_store(ci, 0, memory_order_relaxed, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const atomic_int *' (aka 'const _Atomic(int) *') invalid)}}
   __opencl_atomic_store(i_g, 0, memory_order_relaxed, memory_scope_work_group);
   __opencl_atomic_store(i_l, 0, memory_order_relaxed, memory_scope_work_group);
   __opencl_atomic_store(i_p, 0, memory_order_relaxed, memory_scope_work_group);
   __opencl_atomic_store(i_c, 0, memory_order_relaxed, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-constant _Atomic type ('__constant atomic_int *' (aka '__constant _Atomic(int) *') invalid)}}
 
   __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_load(p, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_load(d, memory_order_seq_cst, memory_scope_work_group);
-  __opencl_atomic_load(ci, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+  __opencl_atomic_load(ci, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const atomic_int *' (aka 'const _Atomic(int) *') invalid)}}
 
   __opencl_atomic_store(i, 1, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_store(p, 1, memory_order_seq_cst, memory_scope_work_group);
   (int)__opencl_atomic_store(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{operand of type 'void' where arithmetic or pointer type is required}}
 
   int exchange_1 = __opencl_atomic_exchange(i, 1, memory_order_seq_cst, memory_scope_work_group);
   int exchange_2 = __opencl_atomic_exchange(I, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to _Atomic}}
 
   __opencl_atomic_fetch_add(i, 1, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_fetch_add(p, 1, memory_order_seq_cst, memory_scope_work_group);
-  __opencl_atomic_fetch_add(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
+  __opencl_atomic_fetch_add(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('atomic_float *' (aka '_Atomic(float) *') invalid)}}
   __opencl_atomic_fetch_and(i, 1, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_fetch_and(p, 1, memory_order_seq_cst, memory_scope_work_group);
-  __opencl_atomic_fetch_and(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to bitwise atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
+  __opencl_atomic_fetch_and(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to bitwise atomic operation must be a pointer to atomic integer ('atomic_float *' (aka '_Atomic(float) *') invalid)}}
 
   __opencl_atomic_fetch_min(i, 1, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_fetch_max(i, 1, memory_order_seq_cst, memory_scope_work_group);
-  __opencl_atomic_fetch_min(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
-  __opencl_atomic_fetch_max(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
+  __opencl_atomic_fetch_min(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('atomic_float *' (aka '_Atomic(float) *') invalid)}}
+  __opencl_atomic_fetch_max(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('atomic_float *' (aka '_Atomic(float) *') invalid)}}
 
   bool cmpexch_1 = __opencl_atomic_compare_exchange_strong(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
   bool cmpexch_2 = __opencl_atomic_compare_exchange_strong(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
-  bool cmpexch_3 = __opencl_atomic_compare_exchange_strong(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *' to parameter of type '__generic float *'}}
-  (void)__opencl_atomic_compare_exchange_strong(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *' to parameter of type '__generic int *' discards qualifiers}}
+  bool cmpexch_3 = __opencl_atomic_compare_exchange_strong(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing 'int *' to parameter of type '__generic float *'}}
+  (void)__opencl_atomic_compare_exchange_strong(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const int *' to parameter of type '__generic int *' discards qualifiers}}
 
   bool cmpexchw_1 = __opencl_atomic_compare_exchange_weak(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
   bool cmpexchw_2 = __opencl_atomic_compare_exchange_weak(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
-  bool cmpexchw_3 = __opencl_atomic_compare_exchange_weak(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *' to parameter of type '__generic float *'}}
-  (void)__opencl_atomic_compare_exchange_weak(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *' to parameter of type '__generic int *' discards qualifiers}}
+  bool cmpexchw_3 = __opencl_atomic_compare_exchange_weak(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing 'int *' to parameter of type '__generic float *'}}
+  (void)__opencl_atomic_compare_exchange_weak(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const int *' to parameter of type '__generic int *' discards qualifiers}}
 
   // Pointers to different address spaces are allowed.
   bool cmpexch_10 = __opencl_atomic_compare_exchange_strong((global atomic_int *)0x308, (constant int *)0x309, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
 
-  __opencl_atomic_init(ci, 0); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
-  __opencl_atomic_store(ci, 0, memory_order_release, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
-  __opencl_atomic_load(ci, memory_order_acquire, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+  __opencl_atomic_init(ci, 0); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const atomic_int *' (aka 'const _Atomic(int) *') invalid)}}
+  __opencl_atomic_store(ci, 0, memory_order_release, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const atomic_int *' (aka 'const _Atomic(int) *') invalid)}}
+  __opencl_atomic_load(ci, memory_order_acquire, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const atomic_int *' (aka 'const _Atomic(int) *') invalid)}}
 
   __opencl_atomic_init(&gn, 456);
-  __opencl_atomic_init(&gn, (void*)0); // expected-warning{{incompatible pointer to integer conversion passing '__generic void *' to parameter of type 'int'}}
+  __opencl_atomic_init(&gn, (void*)0); // expected-warning{{incompatible pointer to integer conversion passing 'void *' to parameter of type 'int'}}
 }
 
 void memory_checks(atomic_int *Ap, int *p, int val) {
   // non-integer memory order argument is casted to integer type.
   (void)__opencl_atomic_load(Ap, 1.0f, memory_scope_work_group);
   float forder;
   (void)__opencl_atomic_load(Ap, forder, memory_scope_work_group);
   struct S s;

test/SemaOpenCL/invalid-block.cl

 // RUN: %clang_cc1 -verify -fblocks -cl-std=CL2.0 %s
 
 // OpenCL v2.0 s6.12.5
 void f0(int (^const bl)());
 // All blocks declarations must be const qualified and initialized.
 void f1() {
   int (^bl1)(void) = ^() {
     return 1;
   };
   int (^const bl2)(void) = ^() {
     return 1;
   };
   f0(bl1);
   f0(bl2);
-  bl1 = bl2;          // expected-error{{invalid operands to binary expression ('int (__generic ^const)(void)' and 'int (__generic ^const)(void)')}}
+  bl1 = bl2;          // expected-error{{invalid operands to binary expression ('int (^const)(void)' and 'int (^const)(void)')}}
   int (^const bl3)(); // expected-error{{invalid block variable declaration - must be initialized}}
 }
 
 // A block with extern storage class is not allowed.
 extern int (^bl)(void) = ^() { // expected-error{{invalid block variable declaration - using 'extern' storage class is disallowed}}
   return 1;
 };
 void f2() {
   extern int (^bl)(void) = ^() { // expected-error{{invalid block variable declaration - using 'extern' storage class is disallowed}}
     return 1;
   };
 }
 
 // A block cannot be the return value of a function.
 typedef int (^bl_t)(void);
-bl_t f3(bl_t bl); // expected-error{{declaring function return value of type 'bl_t' (aka 'int (__generic ^const)(void)') is not allowed}}
+bl_t f3(bl_t bl); // expected-error{{declaring function return value of type 'bl_t' (aka 'int (^const)(void)') is not allowed}}
 
 struct bl_s {
-  int (^bl)(void); // expected-error {{the 'int (__generic ^const)(void)' type cannot be used to declare a structure or union field}}
+  int (^bl)(void); // expected-error {{the 'int (^const)(void)' type cannot be used to declare a structure or union field}}
 };
 
 void f4() {
   __block int a = 10; // expected-error {{the __block storage type is not permitted}}
 }
 
 // A block with variadic argument is not allowed.
 int (^bl)(int, ...) = ^int(int I, ...) { // expected-error {{invalid prototype, variadic arguments are not allowed in OpenCL}} expected-error {{invalid prototype, variadic arguments are not allowed in OpenCL}}
   return 0;
 };
 typedef int (^bl1_t)(int, ...); // expected-error {{invalid prototype, variadic arguments are not allowed in OpenCL}}
 
 // A block can't be used to declare an array
 typedef int (^bl2_t)(int);
 void f5(int i) {
   bl2_t bl1 = ^(int i) {
     return 1;
   };
   bl2_t bl2 = ^(int i) {
     return 2;
   };
-  bl2_t arr[] = {bl1, bl2}; // expected-error {{array of 'bl2_t' (aka 'int (__generic ^const)(int)') type is invalid in OpenCL}}
+  bl2_t arr[] = {bl1, bl2}; // expected-error {{array of 'bl2_t' (aka 'int (^const)(int)') type is invalid in OpenCL}}
   int tmp = i ? bl1(i)      // expected-error {{block type cannot be used as expression in ternary expression in OpenCL}}
               : bl2(i);     // expected-error {{block type cannot be used as expression in ternary expression in OpenCL}}
 }
 // A block pointer type and all pointer operations are disallowed
-void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(int)') is invalid in OpenCL}}
+void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type 'bl2_t' (aka 'int (^const)(int)') is invalid in OpenCL}}
   bl2_t bl = ^(int i) {
     return 1;
   };
-  bl2_t *p; // expected-error {{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(int)') is invalid in OpenCL}}
-  *bl;      // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}}
-  &bl;      // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}}
+  bl2_t *p; // expected-error {{pointer to type 'bl2_t' (aka 'int (^const)(int)') is invalid in OpenCL}}
+  *bl;      // expected-error {{invalid argument type 'bl2_t' (aka 'int (^const)(int)') to unary expression}}
+  &bl;      // expected-error {{invalid argument type 'bl2_t' (aka 'int (^const)(int)') to unary expression}}
 }
 // A block can't reference another block
 kernel void f7() {
   bl2_t bl1 = ^(int i) {
     return 1;
   };
   void (^bl2)(void) = ^{
     int i = bl1(1); // expected-error {{cannot refer to a block inside block}}

test/SemaOpenCL/invalid-pipes-cl2.0.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0
 
 global pipe int gp;            // expected-error {{type '__global read_only pipe int' can only be used as a function parameter in OpenCL}}
 global reserve_id_t rid;          // expected-error {{the '__global reserve_id_t' type cannot be used to declare a program scope variable}}
 
-extern pipe write_only int get_pipe(); // expected-error {{type '__global write_only pipe int ()' can only be used as a function parameter in OpenCL}}
+extern pipe write_only int get_pipe(); // expected-error {{type 'write_only pipe int ()' can only be used as a function parameter in OpenCL}}
 
 kernel void test_invalid_reserved_id(reserve_id_t ID) { // expected-error {{'reserve_id_t' cannot be used as the type of a kernel parameter}}
 }
 
 void test1(pipe int *p) {// expected-error {{pipes packet types cannot be of reference type}}
 }
 void test2(pipe p) {// expected-error {{missing actual type specifier for pipe}}
 }

test/SemaOpenCL/null_literal.cl

 // RUN: %clang_cc1 -verify %s
-// RUN: %clang_cc1 -cl-std=CL2.0 -DCL20 -verify %s
+// RUN: %clang_cc1 -cl-std=CL2.0 -verify %s
 
 #define NULL ((void*)0)
 
 void foo(){
-
-global int* ptr1 = NULL;
-
-global int* ptr2 = (global void*)0;
-
-constant int* ptr3 = NULL;
-
-constant int* ptr4 = (global void*)0; // expected-error{{initializing '__constant int *' with an expression of type '__global void *' changes address space of pointer}}
-
-#ifdef CL20
-// Accept explicitly pointer to generic address space in OpenCL v2.0.
-global int* ptr5 = (generic void*)0;
+  global int *g1 = NULL;
+  global int *g2 = (global void *)0;
+  global int *g3 = (constant void *)0; // expected-error{{initializing '__global int *' with an expression of type '__constant void *' changes address space of pointer}}
+  global int *g4 = (local void *)0; // expected-error{{initializing '__global int *' with an expression of type '__local void *' changes address space of pointer}}
+  global int *g5 = (private void *)0; // expected-error{{initializing '__global int *' with an expression of type '__private void *' changes address space of pointer}}
+
+  constant int *c1 = NULL;
+  constant int *c2 = (global void *)0; // expected-error{{initializing '__constant int *' with an expression of type '__global void *' changes address space of pointer}}
+  constant int *c3 = (constant void *)0;
+  constant int *c4 = (local void *)0; // expected-error{{initializing '__constant int *' with an expression of type '__local void *' changes address space of pointer}}
+  constant int *c5 = (private void *)0; // expected-error{{initializing '__constant int *' with an expression of type '__private void *' changes address space of pointer}}
+
+  local int *l1 = NULL;
+  local int *l2 = (global void *)0; // expected-error{{initializing '__local int *' with an expression of type '__global void *' changes address space of pointer}}
+  local int *l3 = (constant void *)0; // expected-error{{initializing '__local int *' with an expression of type '__constant void *' changes address space of pointer}}
+  local int *l4 = (local void *)0;
+  local int *l5 = (private void *)0; // expected-error{{initializing '__local int *' with an expression of type '__private void *' changes address space of pointer}}
+
+  private int *p1 = NULL;
+  private int *p2 = (global void *)0; // expected-error{{initializing '__private int *' with an expression of type '__global void *' changes address space of pointer}}
+  private int *p3 = (constant void *)0; // expected-error{{initializing '__private int *' with an expression of type '__constant void *' changes address space of pointer}}
+  private int *p4 = (local void *)0; // expected-error{{initializing '__private int *' with an expression of type '__local void *' changes address space of pointer}}
+  private int *p5 = (private void *)0;
+
+#if __OPENCL_C_VERSION__ >= 200
+  // Assigning a pointer to a pointer to narrower address space causes an error unless there is an valid explicit cast.
+  global int *g6 = (generic void *)0; // expected-error{{initializing '__global int *' with an expression of type '__generic void *' changes address space of pointer}}
+  constant int *c6 = (generic void *)0; // expected-error{{initializing '__constant int *' with an expression of type '__generic void *' changes address space of pointer}}
+  local int *l6 = (generic void *)0; // expected-error{{initializing '__local int *' with an expression of type '__generic void *' changes address space of pointer}}
+  private int *p6 = (generic void *)0; // expected-error{{initializing '__private int *' with an expression of type '__generic void *' changes address space of pointer}}
+
+  global int *g7 = (global void*)(generic void *)0;
+  constant int *c7 = (constant void*)(generic void *)0; //expected-error{{casting '__generic void *' to type '__constant void *' changes address space of pointer}}
+  local int *l7 = (local void*)(generic void *)0;
+  private int *p7 = (private void*)(generic void *)0;
+
+  generic int *ge1 = NULL;
+  generic int *ge2 = (global void *)0;
+  generic int *ge3 = (constant void *)0; // expected-error{{initializing '__generic int *' with an expression of type '__constant void *' changes address space of pointer}}
+  generic int *ge4 = (local void *)0;
+  generic int *ge5 = (private void *)0;
+  generic int *ge6 = (generic void *)0;
 #endif
 
-global int* ptr6 = (local void*)0; // expected-error{{initializing '__global int *' with an expression of type '__local void *' changes address space of pointer}}
-
-bool cmp = ptr1 == NULL;
-
-cmp = ptr1 == (local void*)0; // expected-error{{comparison between  ('__global int *' and '__local void *') which are pointers to non-overlapping address spaces}}
-
-cmp = ptr3 == NULL;
-
+  bool cmp;
+  cmp = g1 == NULL;
+  cmp = g1 == (global void *)0;
+  cmp = g1 == (constant void *)0; // expected-error{{comparison between  ('__global int *' and '__constant void *') which are pointers to non-overlapping address spaces}}
+  cmp = g1 == (local void *)0; // expected-error{{comparison between  ('__global int *' and '__local void *') which are pointers to non-overlapping address spaces}}
+  cmp = g1 == (private void *)0; // expected-error{{comparison between  ('__global int *' and '__private void *') which are pointers to non-overlapping address spaces}}
+
+#if __OPENCL_C_VERSION__ >= 200
+  cmp = g1 == (generic void *)0;
+
+  cmp = ge1 == NULL;
+  cmp = ge1 == (global void *)0;
+  cmp = ge1 == (constant void *)0; // expected-error{{comparison between  ('__generic int *' and '__constant void *') which are pointers to non-overlapping address spaces}}
+  cmp = ge1 == (local void *)0;
+  cmp = ge1 == (private void *)0;
+  cmp = ge1 == (generic void *)0;
+#endif
 }

test/SemaOpenCL/vector_conv_invalid.cl

void vector_conv_invalid(const global int4 *const_global_ptr) {
   int4 e = (int4)u; // expected-error{{invalid conversion between ext-vector type 'int4' (vector of 4 'int' values) and 'uint4' (vector of 4 'unsigned int' values)}}
 
   uint3 u4 = (uint3)u; // expected-error{{invalid conversion between ext-vector type 'uint3' (vector of 3 'unsigned int' values) and 'uint4' (vector of 4 'unsigned int' values)}}
 
   e = (const int4)i;
   e = (constant int4)i;
   e = (private int4)i;
 
-  private int4 *private_ptr = (const private int4 *)const_global_ptr; // expected-error{{casting 'const __global int4 *' to type 'const int4 *' changes address space of pointer}}
+  private int4 *private_ptr = (const private int4 *)const_global_ptr; // expected-error{{casting 'const __global int4 *' to type 'const __private int4 *' changes address space of pointer}}
   global int4 *global_ptr = const_global_ptr;                 // expected-warning {{initializing '__global int4 *' with an expression of type 'const __global int4 *' discards qualifiers}}
   global_ptr = (global int4 *)const_global_ptr;
 }

test/SemaTemplate/address_space-dependent.cpp

 
 template <int I>
 void neg() {
   __attribute__((address_space(I))) int *bounds; // expected-error {{address space is negative}}
 }
 
 template <long int I>
 void tooBig() {
-  __attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388598)}}
+  __attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (4194294)}}
 }
 
 template <long int I>
 void correct() {
   __attribute__((address_space(I))) int *bounds;
 }
 
 template <int I, int J>

 
 int main() {  
   int __attribute__((address_space(1))) * p1;
   int p = GetAddressSpaceValue(p1);
 
   car<1, 2, 3>(); // expected-note {{in instantiation of function template specialization 'car<1, 2, 3>' requested here}}
   HasASTemplateFields<1> HASTF;
   neg<-1>(); // expected-note {{in instantiation of function template specialization 'neg<-1>' requested here}}
-  correct<0x7FFFF6>();
+  correct<0x3FFFF6>();
   tooBig<8388650>(); // expected-note {{in instantiation of function template specialization 'tooBig<8388650>' requested here}}
 
   __attribute__((address_space(1))) char *x;
   __attribute__((address_space(2))) char *y;
   cmp<1, 2>(x, y); // expected-note {{in instantiation of function template specialization 'cmp<1, 2>' requested here}}
 
   fooFunction<1> ff;
   ff.get_0(); // expected-note {{in instantiation of member function 'fooFunction<1>::get_0' requested here}}
   ff.qf();
   ff.test3(); // expected-note {{in instantiation of member function 'fooFunction<1>::test3' requested here}}
   
   static_assert(partial_spec_deduce_as<int __attribute__((address_space(3))) *>::value == 3, "address space value has been incorrectly deduced"); 
 
   return 0;
 }