[OpenCL] Fixed address space cast in C style cast of C++ parsing

[Anastasia]

Summary

C style cast in OpenCL C++ was ignoring the address space conversions from OpenCL C and as a result accepting incorrect code to compile. This commit adds correctness checking for address spaces. This makes all OpenCL C address space conversion tests from Sema passing in -cl-std=c++ mode too. Error messages have to be changed to C++ style in multiple places, but semantics of the diagnostics stay the same.

Note: this doesn't fix C++ cast operators. OpenCL C++ s3.4.4 defines special address space cast operators that align with C++'s philosophy of safe casts. The motivation behind this is that address space casts might require address translation (denoted by the addrspacecast IR instruction). In order to avoid operations modifying the type or constness together with translation of addresses, a separate cast operator is introduced. We might still have to fix the semantics of address spaces in static_cast, reinterpret_cast and const_cast but it's currently unclear to me. I will address this as a separate patch.

[rjmccall]

lib/Sema/SemaCast.cpp

2286

Why is this not a method on CastOperation?

2291

Please test the result of getAs instead of separately testing isPointerType.

Why is this check OpenCL-specific? Address spaces are a general language feature.

[Anastasia]

lib/Sema/SemaCast.cpp

2286

Will change!

2291

I think mainly because I am factoring out from the existing OpenCL check. But it probably makes sense that the semantics of this is not different in other languages. I will update it! Thanks!

[Anastasia]

Change AS checking function to be a method of CastOperation.

[Anastasia]

lib/Sema/SemaCast.cpp

2291

After playing with this for a bit longer I discovered that I have to keep the OpenCL check unfortunately.

I found this old commit (d4c5f84/r129583) that says:

C-style casts can add/remove/change address spaces through the reinterpret_cast mechanism.

That's not the same as in OpenCL because it seems for C++ you can cast any AS to any other AS. Therefore, no checks are needed at all. I am not sure if we can come up with a common function for the moment.

The following tests are checking this:

CodeGenCXX/address-space-cast.cpp
SemaCXX/address-space-conversion.cpp

Do you think it would make sense to rename this method with OpenCL-something or keep in case may be CUDA or some other languages might need similar functionality...

[rjmccall]

lib/Sema/SemaCast.cpp

2291

I think you can leave it alone for now, but please add a comment explaining the reasoning as best you see it, and feel free to express uncertainty about the right rule.

Don't take QualType by const &, by the way. It's a cheap-to-copy value type and should always be passed by value.

[Anastasia]

• Changed to pass QualType by value.
• Improved description of new method.
• Added a comment about the OpenCL rules.

[Anastasia]

lib/Sema/SemaCast.cpp

2291

My general understand of the address spaces in C and C++ that they are intended to be more general than in OpenCL (i.e. they don't reflect any particular memory system). Perhaps, it makes sense that we have more restrictions in OpenCL or other similar languages.

[rjmccall]

LGTM.

lib/Sema/SemaCast.cpp

2291

No, their primary use is to reflect underlying memory systems, like near vs. far pointers on 32-on-64 targets or specialized address spaces like gs-segment addressing on x86-64, and it doesn't make sense to allow arbitrary conversions any more than it does for OpenCL. The Embedded C spec has rules about overlapping address spaces, and while it doesn't say that casts between non-overlapping address spaces are ill-formed, it probably ought to. Regardless, if we've permitted arbitrary casts in the past, we'll need to investigate the source compatibility issues before imposing restrictions.

There have also been proposals for "superficial" address spaces that are eliminated during lowering which might have their own restrictions.

[Anastasia]

lib/Sema/SemaCast.cpp

2291

Ok, I see. Even if they seem to cover variety of different memory architectures, it probably doesn't make sense to have arbitrary conversions indeed. I think heterogeneous C++ group was looking at regulating the address spaces. I will try to get in touch with them and investigate.

[Anastasia]

Committed in r344148. The review link was forgotten and therefore didn't get closed automatically.

lib/Sema/SemaCast.cpp

struct CastOperation {
 
     bool isPlaceholder() const {
       return PlaceholderKind != 0;
     }
     bool isPlaceholder(BuiltinType::Kind K) const {
       return PlaceholderKind == K;
     }
 
+    // Language specific cast restrictions for address spaces.
+    void checkAddressSpaceCast(QualType SrcType, QualType DestType);
+
     void checkCastAlign() {
       Self.CheckCastAlign(SrcExpr.get(), DestType, OpRange);
     }
 
     void checkObjCConversion(Sema::CheckedConversionKind CCK) {
       assert(Self.getLangOpts().allowsNonTrivialObjCLifetimeQualifiers());
 
       Expr *src = SrcExpr.get();
static TryCastResult TryReinterpretCast(Sema &Self, ExprResult &SrcExpr,
   // C++ 5.2.10p7: A pointer to an object can be explicitly converted to
   //   a pointer to an object of different type.
   // Void pointers are not specified, but supported by every compiler out there.
   // So we finish by allowing everything that remains - it's got to be two
   // object pointers.
   return SuccessResult;
 }
 
+void CastOperation::checkAddressSpaceCast(QualType SrcType, QualType DestType) {
+  // In OpenCL only conversions between pointers to objects in overlapping
+  // addr spaces are allowed. v2.0 s6.5.5 - Generic addr space overlaps
+  // with any named one, except for constant.
+  if (Self.getLangOpts().OpenCL) {
+    auto SrcPtrType = SrcType->getAs<PointerType>();
+    if (!SrcPtrType)
+      return;
+    auto DestPtrType = DestType->getAs<PointerType>();
+    if (!DestPtrType)
+      return;
+    if (!DestPtrType->isAddressSpaceOverlapping(*SrcPtrType)) {
+      Self.Diag(OpRange.getBegin(),
+                diag::err_typecheck_incompatible_address_space)
+          << SrcType << DestType << Sema::AA_Casting
+          << SrcExpr.get()->getSourceRange();
+      SrcExpr = ExprError();
+    }
+  }
+}
+
 void CastOperation::CheckCXXCStyleCast(bool FunctionalStyle,
                                        bool ListInitialization) {
   assert(Self.getLangOpts().CPlusPlus);
 
   // Handle placeholders.
   if (isPlaceholder()) {
     // C-style casts can resolve __unknown_any types.
     if (claimPlaceholder(BuiltinType::UnknownAny)) {
if (SrcExpr.get()->getType() == Self.Context.OverloadTy) {
         Self.NoteAllOverloadCandidates(SrcExpr.get());
       }
     } else {
       diagnoseBadCast(Self, msg, (FunctionalStyle ? CT_Functional : CT_CStyle),
                       OpRange, SrcExpr.get(), DestType, ListInitialization);
     }
   }
 
+  checkAddressSpaceCast(SrcExpr.get()->getType(), DestType);
+
   if (isValidCast(tcr)) {
     if (Kind == CK_BitCast)
       checkCastAlign();
   } else {
     SrcExpr = ExprError();
   }
 }
 
void CastOperation::CheckCStyleCast() {
   }
   SrcExpr = Self.DefaultFunctionArrayLvalueConversion(SrcExpr.get());
   if (SrcExpr.isInvalid())
     return;
   QualType SrcType = SrcExpr.get()->getType();
 
   assert(!SrcType->isPlaceholderType());
 
-  // OpenCL v1 s6.5: Casting a pointer to address space A to a pointer to
-  // address space B is illegal.
-  if (Self.getLangOpts().OpenCL && DestType->isPointerType() &&
-      SrcType->isPointerType()) {
-    const PointerType *DestPtr = DestType->getAs<PointerType>();
-    if (!DestPtr->isAddressSpaceOverlapping(*SrcType->getAs<PointerType>())) {
-      Self.Diag(OpRange.getBegin(),
-                diag::err_typecheck_incompatible_address_space)
-          << SrcType << DestType << Sema::AA_Casting
-          << SrcExpr.get()->getSourceRange();
-      SrcExpr = ExprError();
-      return;
-    }
-  }
+  checkAddressSpaceCast(SrcType, DestType);
+  if (SrcExpr.isInvalid())
 
   if (Self.RequireCompleteType(OpRange.getBegin(), DestType,
                                diag::err_typecheck_cast_to_incomplete)) {
     SrcExpr = ExprError();
     return;
   }
 
   if (!DestType->isScalarType() && !DestType->isVectorType()) {

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

 // RUN: %clang_cc1 %s -ffake-address-space-map -verify -pedantic -fsyntax-only -DCONSTANT -cl-std=CL2.0
 // RUN: %clang_cc1 %s -ffake-address-space-map -verify -pedantic -fsyntax-only -DGLOBAL -cl-std=CL2.0
 // RUN: %clang_cc1 %s -ffake-address-space-map -verify -pedantic -fsyntax-only -DGENERIC -cl-std=CL2.0
+// RUN: %clang_cc1 %s -ffake-address-space-map -verify -pedantic -fsyntax-only -DCONSTANT -cl-std=c++
+// RUN: %clang_cc1 %s -ffake-address-space-map -verify -pedantic -fsyntax-only -DGLOBAL -cl-std=c++
+// RUN: %clang_cc1 %s -ffake-address-space-map -verify -pedantic -fsyntax-only -DGENERIC -cl-std=c++
 
 /* OpenCLC v2.0 adds a set of restrictions for conversions between pointers to
 *  different address spaces, mainly described in Sections 6.5.5 and 6.5.6.
 *
 *  It adds notion of overlapping address spaces. The main differention is that
 *  an unnamed address space is added, called '__generic'. Pointers to the
 *  generic address space can be interchangabley used with pointers to any
 *  other address space except for __constant address space (Section 6.5.5).
 *
 *  Based on this there are 3 sets of tests: __generic, named (__global in this
 *  case), and __constant, that should cover all program paths for CL address
 *  space conversions used in initialisations, assignments, casts, comparisons
 *  and arithmetic operations.
 */
 
 #ifdef GENERIC
-#define AS generic
-#define AS_COMP local
-#define AS_INCOMP constant
+#define AS __generic
+#define AS_COMP __local
+#define AS_INCOMP __constant
 #endif
 
 #ifdef GLOBAL
-#define AS global
-#define AS_COMP global
-#define AS_INCOMP local
+#define AS __global
+#define AS_COMP __global
+#define AS_INCOMP __local
 #endif
 
 #ifdef CONSTANT
-#define AS constant
-#define AS_COMP constant
-#define AS_INCOMP global
+#define AS __constant
+#define AS_COMP __constant
+#define AS_INCOMP __global
 #endif
 
-void f_glob(global int *arg_glob) {}
+void f_glob(__global int *arg_glob) {}
 #ifndef GLOBAL
-// expected-note@-2{{passing argument to parameter 'arg_glob' here}}
+#if !__OPENCL_CPP_VERSION__
+// expected-note@-3{{passing argument to parameter 'arg_glob' here}}
+#else
+// expected-note-re@-5{{candidate function not viable: address space mismatch in 1st argument ('__{{generic|constant}} int *'), parameter type must be '__global int *'}}
+#endif
 #endif
 
-void f_loc(local int *arg_loc) {
-} // expected-note@-1{{passing argument to parameter 'arg_loc' here}}
+void f_loc(__local int *arg_loc) {}
+#if !__OPENCL_CPP_VERSION__
+// expected-note@-2{{passing argument to parameter 'arg_loc' here}}
+#else
+// expected-note-re@-4{{candidate function not viable: address space mismatch in 1st argument ('__{{global|generic|constant}} int *'), parameter type must be '__local int *'}}
+#endif
 
-void f_const(constant int *arg_const) {}
+void f_const(__constant int *arg_const) {}
 #ifndef CONSTANT
-// expected-note@-2{{passing argument to parameter 'arg_const' here}}
+#if !__OPENCL_CPP_VERSION__
+// expected-note@-3{{passing argument to parameter 'arg_const' here}}
+#else
+// expected-note-re@-5{{candidate function not viable: address space mismatch in 1st argument ('__{{global|generic}} int *'), parameter type must be '__constant int *'}}
+#endif
 #endif
 
-void f_priv(private int *arg_priv) {
-} // expected-note@-1{{passing argument to parameter 'arg_priv' here}}
+void f_priv(__private int *arg_priv) {}
+#if !__OPENCL_CPP_VERSION__
+// expected-note@-2{{passing argument to parameter 'arg_priv' here}}
+#else
+// expected-note-re@-4{{candidate function not viable: address space mismatch in 1st argument ('__{{global|generic|constant}} int *'), parameter type must be 'int *'}}
+#endif
 
-void f_gen(generic int *arg_gen) {}
+void f_gen(__generic int *arg_gen) {}
 #ifdef CONSTANT
-// expected-note@-2{{passing argument to parameter 'arg_gen' here}}
+#if !__OPENCL_CPP_VERSION__
+// expected-note@-3{{passing argument to parameter 'arg_gen' here}}
+#else
+// expected-note@-5{{candidate function not viable: address space mismatch in 1st argument ('__constant int *'), parameter type must be '__generic int *'}}
+#endif
 #endif
 
-void test_conversion(global int *arg_glob, local int *arg_loc,
-                     constant int *arg_const, private int *arg_priv,
-                     generic int *arg_gen) {
+void test_conversion(__global int *arg_glob, __local int *arg_loc,
+                     __constant int *arg_const, __private int *arg_priv,
+                     __generic int *arg_gen) {
 
   AS int *var_init1 = arg_glob;
 #ifdef CONSTANT
-// expected-error@-2{{initializing '__constant int *' with an expression of type '__global int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{initializing '__constant int *' with an expression of type '__global int *' changes address space of pointer}}
+#else
+// expected-error@-5{{cannot initialize a variable of type '__constant int *' with an lvalue of type '__global int *'}}
+#endif
 #endif
 
   AS int *var_init2 = arg_loc;
 #ifndef GENERIC
-// expected-error-re@-2{{initializing '__{{global|constant}} int *' with an expression of type '__local int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{initializing '__{{global|constant}} int *' with an expression of type '__local int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{cannot initialize a variable of type '__{{global|constant}} int *' with an lvalue of type '__local int *'}}
+#endif
 #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}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{initializing '__{{global|generic}} int *' with an expression of type '__constant int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{cannot initialize a variable of type '__{{global|generic}} int *' with an lvalue of type '__constant int *'}}
+#endif
 #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}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{initializing '__{{global|constant}} int *' with an expression of type 'int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{cannot initialize a variable of type '__{{global|constant}} int *' with an lvalue of type 'int *'}}
+#endif
 #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}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{initializing '__{{global|constant}} int *' with an expression of type '__generic int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{cannot initialize a variable of type '__{{global|constant}} int *' with an lvalue of type '__generic int *'}}
+#endif
 #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;
#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;
   var_impl = arg_glob;
 #ifdef CONSTANT
-// expected-error@-2{{assigning '__global int *' to '__constant int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{assigning '__global int *' to '__constant int *' changes address space of pointer}}
+#else
+// expected-error@-5{{assigning to '__constant int *' from incompatible type '__global int *'}}
+#endif
 #endif
 
   var_impl = arg_loc;
 #ifndef GENERIC
-// expected-error-re@-2{{assigning '__local int *' to '__{{global|constant}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{assigning '__local int *' to '__{{global|constant}} int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type '__local int *'}}
+#endif
 #endif
 
   var_impl = arg_const;
 #ifndef CONSTANT
-// expected-error-re@-2{{assigning '__constant int *' to '__{{global|generic}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{assigning '__constant int *' to '__{{global|generic}} int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{assigning to '__{{global|generic}} int *' from incompatible type '__constant int *'}}
+#endif
 #endif
 
   var_impl = arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{assigning 'int *' to '__{{global|constant}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{assigning 'int *' to '__{{global|constant}} int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type 'int *'}}
+#endif
 #endif
 
   var_impl = arg_gen;
 #ifndef GENERIC
-// expected-error-re@-2{{assigning '__generic int *' to '__{{global|constant}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{assigning '__generic int *' to '__{{global|constant}} int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type '__generic int *'}}
+#endif
 #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;
#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;
   int b = var_cmp != arg_glob;
 #ifdef CONSTANT
-// expected-error@-2{{comparison between  ('__constant int *' and '__global int *') which are pointers to non-overlapping address spaces}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{comparison between  ('__constant int *' and '__global int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error@-5{{comparison of distinct pointer types ('__constant int *' and '__global int *')}}
+#endif
 #endif
 
   b = var_cmp != arg_loc;
 #ifndef GENERIC
-// expected-error-re@-2{{comparison between  ('__{{global|constant}} int *' and '__local int *') which are pointers to non-overlapping address spaces}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{comparison between  ('__{{global|constant}} int *' and '__local int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error-re@-5{{comparison of distinct pointer types ('__{{global|constant}} int *' and '__local int *')}}
+#endif
 #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}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{comparison between  ('__{{global|generic}} int *' and '__constant int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error-re@-5{{comparison of distinct pointer types ('__{{global|generic}} int *' and '__constant int *')}}
+#endif
 #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}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{comparison between  ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error-re@-5{{comparison of distinct pointer types ('__{{global|constant}} int *' and 'int *')}}
+#endif
 #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}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{comparison between  ('__constant int *' and '__generic int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error@-5{{comparison of distinct pointer types ('__constant int *' and '__generic int *')}}
+#endif
 #endif
 
   AS int *var_sub;
   b = var_sub - arg_glob;
 #ifdef CONSTANT
 // expected-error@-2{{arithmetic operation with operands of type  ('__constant int *' and '__global int *') which are pointers to non-overlapping address spaces}}
 #endif
 
#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}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{passing '__{{constant|generic}} int *' to parameter of type '__global int *' changes address space of pointer}}
+#else
+// expected-error@-5{{no matching function for call to 'f_glob'}}
+#endif
 #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_loc(var_sub);
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-2{{passing '__{{global|constant|generic}} int *' to parameter of type '__local int *' changes address space of pointer}}
+#else
+// expected-error@-4{{no matching function for call to 'f_loc'}}
+#endif
 
   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}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{passing '__{{global|generic}} int *' to parameter of type '__constant int *' changes address space of pointer}}
+#else
+// expected-error@-5{{no matching function for call to 'f_const'}}
+#endif
 #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);
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-2{{passing '__{{global|constant|generic}} int *' to parameter of type 'int *' changes address space of pointer}}
+#else
+// expected-error@-4{{no matching function for call to 'f_priv'}}
+#endif
 
   f_gen(var_sub);
 #ifdef CONSTANT
-// expected-error@-2{{passing '__constant int *' to parameter of type '__generic int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{passing '__constant int *' to parameter of type '__generic int *' changes address space of pointer}}
+#else
+// expected-error@-5{{no matching function for call to 'f_gen'}}
+#endif
 #endif
 }
 
 void test_ternary() {
   AS int *var_cond;
-  generic int *var_gen;
-  global int *var_glob;
+  __generic int *var_gen;
+  __global int *var_glob;
   var_gen = 0 ? var_cond : var_glob;
 #ifdef CONSTANT
-// expected-error@-2{{conditional operator with the second and third operands of type  ('__constant int *' and '__global int *') which are pointers to non-overlapping address spaces}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{conditional operator with the second and third operands of type  ('__constant int *' and '__global int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error@-5{{incompatible operand types ('__constant int *' and '__global int *')}}
+#endif
 #endif
 
-  local int *var_loc;
+  __local int *var_loc;
   var_gen = 0 ? var_cond : var_loc;
 #ifndef GENERIC
-// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and '__local int *') which are pointers to non-overlapping address spaces}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and '__local int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error-re@-5{{incompatible operand types ('__{{global|constant}} int *' and '__local int *')}}
+#endif
 #endif
 
-  constant int *var_const;
+  __constant int *var_const;
   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}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{conditional operator with the second and third operands of type  ('__{{global|generic}} int *' and '__constant int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error-re@-5{{incompatible operand types ('__{{global|generic}} int *' and '__constant int *')}}
+#endif
 #endif
 
-  private int *var_priv;
+  __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}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error-re@-5{{incompatible operand types ('__{{global|constant}} int *' and 'int *')}}
+#endif
 #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}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{conditional operator with the second and third operands of type  ('__constant int *' and '__generic int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error@-5{{incompatible operand types ('__constant int *' and '__generic int *')}}
+#endif
 #endif
 
   void *var_void_gen;
-  global char *var_glob_ch;
+  __global char *var_glob_ch;
   var_void_gen = 0 ? var_cond : var_glob_ch;
+#if __OPENCL_CPP_VERSION__
+// expected-error-re@-2{{incompatible operand types ('__{{constant|global|generic}} int *' and '__global char *')}}
+#else
 #ifdef CONSTANT
-// expected-error@-2{{conditional operator with the second and third operands of type  ('__constant int *' and '__global char *') which are pointers to non-overlapping address spaces}}
+// expected-error@-5{{conditional operator with the second and third operands of type  ('__constant int *' and '__global char *') which are pointers to non-overlapping address spaces}}
 #else
-// expected-warning-re@-4{{pointer type mismatch ('__{{global|generic}} int *' and '__global char *')}}
+// expected-warning-re@-7{{pointer type mismatch ('__{{global|generic}} int *' and '__global char *')}}
+#endif
 #endif
 
-  local char *var_loc_ch;
+  __local char *var_loc_ch;
   var_void_gen = 0 ? var_cond : var_loc_ch;
+#if __OPENCL_CPP_VERSION__
+// expected-error-re@-2{{incompatible operand types ('__{{constant|global|generic}} int *' and '__local char *')}}
+#else
 #ifndef GENERIC
-// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and '__local char *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-5{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and '__local char *') which are pointers to non-overlapping address spaces}}
 #else
-// expected-warning@-4{{pointer type mismatch ('__generic int *' and '__local char *')}}
+// expected-warning@-7{{pointer type mismatch ('__generic int *' and '__local char *')}}
+#endif
 #endif
 
-  constant void *var_void_const;
-  constant char *var_const_ch;
+  __constant void *var_void_const;
+  __constant char *var_const_ch;
   var_void_const = 0 ? var_cond : var_const_ch;
+#if __OPENCL_CPP_VERSION__
+// expected-error-re@-2{{incompatible operand types ('__{{constant|global|generic}} int *' and '__constant char *')}}
+#else
 #ifndef CONSTANT
-// 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}}
+// expected-error-re@-5{{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 *')}}
+// expected-warning@-7{{pointer type mismatch ('__constant int *' and '__constant char *')}}
+#endif
 #endif
 
-  private char *var_priv_ch;
+  __private char *var_priv_ch;
   var_void_gen = 0 ? var_cond : var_priv_ch;
+#if __OPENCL_CPP_VERSION__
+// expected-error-re@-2{{incompatible operand types ('__{{constant|global|generic}} int *' and 'char *')}}
+#else
 #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@-5{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and 'char *') which are pointers to non-overlapping address spaces}}
 #else
-// expected-warning@-4{{pointer type mismatch ('__generic int *' and 'char *')}}
+// expected-warning@-7{{pointer type mismatch ('__generic int *' and 'char *')}}
+#endif
 #endif
 
-  generic char *var_gen_ch;
+  __generic char *var_gen_ch;
   var_void_gen = 0 ? var_cond : var_gen_ch;
+#if __OPENCL_CPP_VERSION__
+// expected-error-re@-2{{incompatible operand types ('__{{constant|global|generic}} int *' and '__generic char *')}}
+#else
 #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}}
+// expected-error@-5{{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 *')}}
+// expected-warning-re@-7{{pointer type mismatch ('__{{global|generic}} int *' and '__generic char *')}}
+#endif
 #endif
 }
 
 void test_pointer_chains() {
   AS int *AS *var_as_as_int;
   AS int *AS_COMP *var_asc_as_int;
   AS_INCOMP int *AS_COMP *var_asc_asn_int;
   AS_COMP int *AS_COMP *var_asc_asc_int;
 
   // Case 1:
   //  * address spaces of corresponded most outer pointees overlaps, their canonical types are equal
   //  * CVR, address spaces and canonical types of the rest of pointees are equivalent.
   var_as_as_int = 0 ? var_as_as_int : var_asc_as_int;
-
+#if __OPENCL_CPP_VERSION__
+#ifdef GENERIC
+// expected-error@-3{{incompatible operand types ('__generic int *__generic *' and '__generic int *__local *')}}
+#endif
+#endif
   // Case 2: Corresponded inner pointees has non-overlapping address spaces.
   var_as_as_int = 0 ? var_as_as_int : var_asc_asn_int;
-// expected-warning-re@-1{{pointer type mismatch ('__{{(generic|global|constant)}} int *__{{(generic|global|constant)}} *' and '__{{(local|global|constant)}} int *__{{(constant|local|global)}} *')}}
+#if !__OPENCL_CPP_VERSION__
+// expected-warning-re@-2{{pointer type mismatch ('__{{(generic|global|constant)}} int *__{{(generic|global|constant)}} *' and '__{{(local|global|constant)}} int *__{{(constant|local|global)}} *')}}
+#else
+// expected-error-re@-4{{incompatible operand types ('__{{(generic|global|constant)}} int *__{{(generic|global|constant)}} *' and '__{{(local|global|constant)}} int *__{{(constant|local|global)}} *')}}
+#endif
 
   // Case 3: Corresponded inner pointees has overlapping but not equivalent address spaces.
 #ifdef GENERIC
   var_as_as_int = 0 ? var_as_as_int : var_asc_asc_int;
-// expected-warning-re@-1{{pointer type mismatch ('__{{(generic|global|constant)}} int *__{{(generic|global|constant)}} *' and '__{{(local|global|constant)}} int *__{{(local|global|constant)}} *')}}
+#if !__OPENCL_CPP_VERSION__
+// expected-warning-re@-2{{pointer type mismatch ('__{{(generic|global|constant)}} int *__{{(generic|global|constant)}} *' and '__{{(local|global|constant)}} int *__{{(local|global|constant)}} *')}}
+#else
+// expected-error-re@-4{{incompatible operand types ('__{{generic|global|constant}} int *__{{generic|global|constant}} *' and '__{{local|global|constant}} int *__{{local|global|constant}} *')}}
+#endif
 #endif
 }

test/SemaOpenCL/address-spaces.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
 // RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only
 // RUN: %clang_cc1 %s -cl-std=c++ -verify -pedantic -fsyntax-only
 
 __constant int ci = 1;
 
 __kernel void foo(__global int *gip) {
   __local int li;
   __local int lj = 2; // expected-error {{'__local' variable cannot have an initializer}}
 
   int *ip;
-// FIXME: Temporarily disable part of the test that doesn't work for C++ yet.
-#if !__OPENCL_CPP_VERSION__
-#if __OPENCL_C_VERSION__ < 200
+#if ((!__OPENCL_CPP_VERSION__) && (__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;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{assigning '__constant int *' to '__generic int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{assigning to '__generic int *' from incompatible type '__constant int *'}}
+#endif
 #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}}
-
-  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}}
-
-  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}}
-
-  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}}
+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}}
+
+  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}}
+
+  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}}
+
+  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}}
 }
 
-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;
+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;
 }
 
 __private int func_return_priv(void);       //expected-error {{return value cannot be qualified with address space}}
 __global int func_return_global(void);      //expected-error {{return value cannot be qualified with address space}}
 __local int func_return_local(void);        //expected-error {{return value cannot be qualified with address space}}
 __constant int func_return_constant(void);  //expected-error {{return value cannot be qualified with address space}}
 #if __OPENCL_C_VERSION__ >= 200
 __generic int func_return_generic(void);    //expected-error {{return value cannot be qualified with address space}}
 #endif
 
 void func_multiple_addr(void) {
   typedef __private int private_int_t;
   __private __local int var1;   // expected-error {{multiple address spaces specified for type}}
   __private __local int *var2;  // expected-error {{multiple address spaces specified for type}}
   __local private_int_t var3;   // expected-error {{multiple address spaces specified for type}}
   __local private_int_t *var4;  // expected-error {{multiple address spaces specified for type}}
   __private private_int_t var5; // expected-warning {{multiple identical address spaces specified for type}}
   __private private_int_t *var6;// expected-warning {{multiple identical address spaces specified for type}}
-#endif // !__OPENCL_CXX_VERSION__
 }