[OpenCL] Pretty print __private addr space

[Anastasia]

Summary

Historically Clang didn't have representation of __private and therefore it has never been printed in types. It appeared to be the same as no address space and appeared to be confusing to the users. Also it was unnecessarily exposing the implementation details of Clang.

We have added a representation of it in AST https://reviews.llvm.org/D35082 but it hasn't been added to the type printing yet.

Current patch adds printing of __private in type printer. The main motivation for this are:

1. Avoid situations where __private addr space is missing in the type of diagnostics even if it was provided in the source code.
2. Improve frontend debugging support because we will now be able to differentiate between __private and absence of addr space in AST dumps.

There is one potential issue that some diagnostics will become more verbose with extra __private addr spaces being printed that were deduced implicitly. This aligns however with the rest of behaviour i.e. deduced __generic addr space is currently printed in AST dump and diagnostics. It is also reasonable for deduced addr spaces to appear in type printing because it can be explained by the language rules. However, absence of __private from type printing has been an implementation detail of Clang that never was intuitive to the users.

[bader]

LGTM, but I'd like someone who works on OpenCL front-end to approve.
+@AlexeySotkin

[svenvh]

LGTM.

[AlexeySotkin]

clang/test/SemaOpenCL/access-qualifier.cl

28

Minor. An error message like this looks a bit confusing to me. User might wonder whether parameters are incompatible because of address space or because of access qualifiers. Should it print passing '__private img1d_wo' (aka '__private __write_only image1d_t') to parameter of incompatible type '__private __read_only image1d_t'? In this case it is clear that there is a mismatch in access qualifiers.

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

373

Again, I think it would be more clear if it was: passing '__constant int *__private' to parameter of type '__generic int *__private' changes address space of pointer

[Anastasia]

clang/test/SemaOpenCL/access-qualifier.cl

28

Yes, I agree. However, we are printing the full QualType in both places in this diagnostics, so the problem is that the addr space is either not deduced or being dropped from QualType somewhere. I don't think we have got addr spaces working yet for all cases correctly and printing __private have revealed a number of such issues. I suggest however to fix them in isolation case by case as we discover them. This commit is already pretty big and I don't want to expand it even more. I have opened a bug to track this issues: https://bugs.llvm.org/show_bug.cgi?id=44294
Hopefully we can fix it asap. Does it make sense?

[AlexeySotkin]

clang/test/SemaOpenCL/access-qualifier.cl

28

Yes, fair enough.

clang/lib/AST/TypePrinter.cpp

std::string Qualifiers::getAddrSpaceAsString(LangAS AS) {
   switch (AS) {
   case LangAS::Default:
     return "";
   case LangAS::opencl_global:
     return "__global";
   case LangAS::opencl_local:
     return "__local";
   case LangAS::opencl_private:
-    return "";
+    return "__private";
   case LangAS::opencl_constant:
     return "__constant";
   case LangAS::opencl_generic:
     return "__generic";
   case LangAS::cuda_device:
     return "__device__";
   case LangAS::cuda_constant:
     return "__constant__";

clang/lib/Sema/SemaDecl.cpp

static bool isOpenCLSizeDependentType(ASTContext &C, QualType Ty) {
   // integers other than by their names.
   StringRef SizeTypeNames[] = {"size_t", "intptr_t", "uintptr_t", "ptrdiff_t"};
 
   // Remove typedefs one by one until we reach a typedef
   // for a size dependent type.
   QualType DesugaredTy = Ty;
   do {
     ArrayRef<StringRef> Names(SizeTypeNames);
-    auto Match = llvm::find(Names, DesugaredTy.getAsString());
+    auto Match = llvm::find(Names, DesugaredTy.getUnqualifiedType().getAsString());
     if (Names.end() != Match)
       return true;
 
     Ty = DesugaredTy;
     DesugaredTy = Ty.getSingleStepDesugaredType(C);
   } while (DesugaredTy != Ty);
 
   return false;

clang/test/AST/language_address_space_attribute.cpp

void langas() {
   [[clang::opencl_local]] int *z_local;
 
   // CHECK: VarDecl {{.*}} x_constant '__constant int *'
   __attribute__((opencl_constant)) int *x_constant;
 
   // CHECK: VarDecl {{.*}} z_constant '__constant int *'
   [[clang::opencl_constant]] int *z_constant;
 
-  // CHECK: VarDecl {{.*}} x_private 'int *'
+  // CHECK: VarDecl {{.*}} x_private '__private int *'
   __attribute__((opencl_private)) int *x_private;
 
-  // CHECK: VarDecl {{.*}} z_private 'int *'
+  // CHECK: VarDecl {{.*}} z_private '__private int *'
   [[clang::opencl_private]] int *z_private;
 
   // CHECK: VarDecl {{.*}} x_generic '__generic int *'
   __attribute__((opencl_generic)) int *x_generic;
 
   // CHECK: VarDecl {{.*}} z_generic '__generic int *'
   [[clang::opencl_generic]] int *z_generic;
 }

clang/test/Index/opencl-types.cl

void kernel testFloatTypes() {
   half scalarHalf;
   half4 vectorHalf;
   float scalarFloat;
   float4 vectorFloat;
   double scalarDouble;
   double4 vectorDouble;
 }
 
-// CHECK: VarDecl=scalarHalf:11:8 (Definition){{( \(invalid\))?}} [type=half] [typekind=Half] [isPOD=1]
-// CHECK: VarDecl=vectorHalf:12:9 (Definition) [type=half4] [typekind=Typedef] [canonicaltype=half __attribute__((ext_vector_type(4)))] [canonicaltypekind=ExtVector] [isPOD=1]
-// CHECK: VarDecl=scalarFloat:13:9 (Definition) [type=float] [typekind=Float] [isPOD=1]
-// CHECK: VarDecl=vectorFloat:14:10 (Definition) [type=float4] [typekind=Typedef] [canonicaltype=float __attribute__((ext_vector_type(4)))] [canonicaltypekind=ExtVector] [isPOD=1]
-// CHECK: VarDecl=scalarDouble:15:10 (Definition){{( \(invalid\))?}} [type=double] [typekind=Double] [isPOD=1]
-// CHECK: VarDecl=vectorDouble:16:11 (Definition){{( \(invalid\))?}} [type=double4] [typekind=Typedef] [canonicaltype=double __attribute__((ext_vector_type(4)))] [canonicaltypekind=ExtVector] [isPOD=1]
+// CHECK: VarDecl=scalarHalf:11:8 (Definition){{( \(invalid\))?}} [type=__private half] [typekind=Half] [isPOD=1]
+// CHECK: VarDecl=vectorHalf:12:9 (Definition) [type=__private half4] [typekind=Typedef] [canonicaltype=half __private __attribute__((ext_vector_type(4)))] [canonicaltypekind=ExtVector] [isPOD=1]
+// CHECK: VarDecl=scalarFloat:13:9 (Definition) [type=__private float] [typekind=Float] [isPOD=1]
+// CHECK: VarDecl=vectorFloat:14:10 (Definition) [type=__private float4] [typekind=Typedef] [canonicaltype=float __private __attribute__((ext_vector_type(4)))] [canonicaltypekind=ExtVector] [isPOD=1]
+// CHECK: VarDecl=scalarDouble:15:10 (Definition){{( \(invalid\))?}} [type=__private double] [typekind=Double] [isPOD=1]
+// CHECK: VarDecl=vectorDouble:16:11 (Definition){{( \(invalid\))?}} [type=__private double4] [typekind=Typedef] [canonicaltype=double __private __attribute__((ext_vector_type(4)))] [canonicaltypekind=ExtVector] [isPOD=1]
 
 #pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable
 
 void kernel OCLImage1dROTest(read_only image1d_t scalarOCLImage1dRO);
 void kernel OCLImage1dArrayROTest(read_only image1d_array_t scalarOCLImage1dArrayRO);
 void kernel OCLImage1dBufferROTest(read_only image1d_buffer_t scalarOCLImage1dBufferRO);
 void kernel OCLImage2dROTest(read_only image2d_t scalarOCLImage2dRO);
 void kernel OCLImage2dArrayROTest(read_only image2d_array_t scalarOCLImage2dArrayRO);
 void kernel OCLImage2dDepthROTest(read_only image2d_depth_t scalarOCLImage2dDepthRO);
 void kernel OCLImage2dArrayDepthROTest(read_only image2d_array_depth_t scalarOCLImage2dArrayDepthRO);
 void kernel OCLImage2dMSAAROTest(read_only image2d_msaa_t scalarOCLImage2dMSAARO);
 void kernel OCLImage2dArrayMSAAROTest(read_only image2d_array_msaa_t scalarOCLImage2dArrayMSAARO);
 void kernel OCLImage2dMSAADepthROTest(read_only image2d_msaa_depth_t scalarOCLImage2dMSAADepthRO);
 void kernel OCLImage2dArrayMSAADepthROTest(read_only image2d_array_msaa_depth_t scalarOCLImage2dArrayMSAADepthRO);
 void kernel OCLImage3dROTest(read_only image3d_t scalarOCLImage3dRO);
 
-// CHECK: ParmDecl=scalarOCLImage1dRO:28:50 (Definition) [type=__read_only image1d_t] [typekind=OCLImage1dRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage1dArrayRO:29:61 (Definition) [type=__read_only image1d_array_t] [typekind=OCLImage1dArrayRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage1dBufferRO:30:63 (Definition) [type=__read_only image1d_buffer_t] [typekind=OCLImage1dBufferRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dRO:31:50 (Definition) [type=__read_only image2d_t] [typekind=OCLImage2dRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayRO:32:61 (Definition) [type=__read_only image2d_array_t] [typekind=OCLImage2dArrayRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dDepthRO:33:61 (Definition) [type=__read_only image2d_depth_t] [typekind=OCLImage2dDepthRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayDepthRO:34:72 (Definition) [type=__read_only image2d_array_depth_t] [typekind=OCLImage2dArrayDepthRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dMSAARO:35:59 (Definition){{( \(invalid\))?}} [type=__read_only image2d_msaa_t] [typekind=OCLImage2dMSAARO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayMSAARO:36:70 (Definition){{( \(invalid\))?}} [type=__read_only image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAARO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dMSAADepthRO:37:70 (Definition){{( \(invalid\))?}} [type=__read_only image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthRO:38:81 (Definition){{( \(invalid\))?}} [type=__read_only image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthRO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage3dRO:39:50 (Definition) [type=__read_only image3d_t] [typekind=OCLImage3dRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dRO:28:50 (Definition) [type=__private __read_only image1d_t] [typekind=OCLImage1dRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dArrayRO:29:61 (Definition) [type=__private __read_only image1d_array_t] [typekind=OCLImage1dArrayRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dBufferRO:30:63 (Definition) [type=__private __read_only image1d_buffer_t] [typekind=OCLImage1dBufferRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dRO:31:50 (Definition) [type=__private __read_only image2d_t] [typekind=OCLImage2dRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayRO:32:61 (Definition) [type=__private __read_only image2d_array_t] [typekind=OCLImage2dArrayRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dDepthRO:33:61 (Definition) [type=__private __read_only image2d_depth_t] [typekind=OCLImage2dDepthRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayDepthRO:34:72 (Definition) [type=__private __read_only image2d_array_depth_t] [typekind=OCLImage2dArrayDepthRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dMSAARO:35:59 (Definition){{( \(invalid\))?}} [type=__private __read_only image2d_msaa_t] [typekind=OCLImage2dMSAARO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayMSAARO:36:70 (Definition){{( \(invalid\))?}} [type=__private __read_only image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAARO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dMSAADepthRO:37:70 (Definition){{( \(invalid\))?}} [type=__private __read_only image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthRO:38:81 (Definition){{( \(invalid\))?}} [type=__private __read_only image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthRO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage3dRO:39:50 (Definition) [type=__private __read_only image3d_t] [typekind=OCLImage3dRO] [isPOD=1]
 
 void kernel OCLImage1dWOTest(write_only image1d_t scalarOCLImage1dWO);
 void kernel OCLImage1dArrayWOTest(write_only image1d_array_t scalarOCLImage1dArrayWO);
 void kernel OCLImage1dBufferWOTest(write_only image1d_buffer_t scalarOCLImage1dBufferWO);
 void kernel OCLImage2dWOTest(write_only image2d_t scalarOCLImage2dWO);
 void kernel OCLImage2dArrayWOTest(write_only image2d_array_t scalarOCLImage2dArrayWO);
 void kernel OCLImage2dDepthWOTest(write_only image2d_depth_t scalarOCLImage2dDepthWO);
 void kernel OCLImage2dArrayDepthWOTest(write_only image2d_array_depth_t scalarOCLImage2dArrayDepthWO);
 void kernel OCLImage2dMSAAWOTest(write_only image2d_msaa_t scalarOCLImage2dMSAAWO);
 void kernel OCLImage2dArrayMSAAWOTest(write_only image2d_array_msaa_t scalarOCLImage2dArrayMSAAWO);
 void kernel OCLImage2dMSAADepthWOTest(write_only image2d_msaa_depth_t scalarOCLImage2dMSAADepthWO);
 void kernel OCLImage2dArrayMSAADepthWOTest(write_only image2d_array_msaa_depth_t scalarOCLImage2dArrayMSAADepthWO);
 void kernel OCLImage3dWOTest(write_only image3d_t scalarOCLImage3dWO);
 
-// CHECK: ParmDecl=scalarOCLImage1dWO:54:51 (Definition) [type=__write_only image1d_t] [typekind=OCLImage1dWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage1dArrayWO:55:62 (Definition) [type=__write_only image1d_array_t] [typekind=OCLImage1dArrayWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage1dBufferWO:56:64 (Definition) [type=__write_only image1d_buffer_t] [typekind=OCLImage1dBufferWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dWO:57:51 (Definition) [type=__write_only image2d_t] [typekind=OCLImage2dWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayWO:58:62 (Definition) [type=__write_only image2d_array_t] [typekind=OCLImage2dArrayWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dDepthWO:59:62 (Definition) [type=__write_only image2d_depth_t] [typekind=OCLImage2dDepthWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayDepthWO:60:73 (Definition) [type=__write_only image2d_array_depth_t] [typekind=OCLImage2dArrayDepthWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dMSAAWO:61:60 (Definition){{( \(invalid\))?}} [type=__write_only image2d_msaa_t] [typekind=OCLImage2dMSAAWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayMSAAWO:62:71 (Definition){{( \(invalid\))?}} [type=__write_only image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAAWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dMSAADepthWO:63:71 (Definition){{( \(invalid\))?}} [type=__write_only image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthWO:64:82 (Definition){{( \(invalid\))?}} [type=__write_only image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthWO] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage3dWO:65:51 (Definition){{( \(invalid\))?}} [type=__write_only image3d_t] [typekind=OCLImage3dWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dWO:54:51 (Definition) [type=__private __write_only image1d_t] [typekind=OCLImage1dWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dArrayWO:55:62 (Definition) [type=__private __write_only image1d_array_t] [typekind=OCLImage1dArrayWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dBufferWO:56:64 (Definition) [type=__private __write_only image1d_buffer_t] [typekind=OCLImage1dBufferWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dWO:57:51 (Definition) [type=__private __write_only image2d_t] [typekind=OCLImage2dWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayWO:58:62 (Definition) [type=__private __write_only image2d_array_t] [typekind=OCLImage2dArrayWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dDepthWO:59:62 (Definition) [type=__private __write_only image2d_depth_t] [typekind=OCLImage2dDepthWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayDepthWO:60:73 (Definition) [type=__private __write_only image2d_array_depth_t] [typekind=OCLImage2dArrayDepthWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dMSAAWO:61:60 (Definition){{( \(invalid\))?}} [type=__private __write_only image2d_msaa_t] [typekind=OCLImage2dMSAAWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayMSAAWO:62:71 (Definition){{( \(invalid\))?}} [type=__private __write_only image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAAWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dMSAADepthWO:63:71 (Definition){{( \(invalid\))?}} [type=__private __write_only image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthWO:64:82 (Definition){{( \(invalid\))?}} [type=__private __write_only image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthWO] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage3dWO:65:51 (Definition){{( \(invalid\))?}} [type=__private __write_only image3d_t] [typekind=OCLImage3dWO] [isPOD=1]
 
 void kernel OCLImage1dRWTest(read_write image1d_t scalarOCLImage1dRW);
 void kernel OCLImage1dArrayRWTest(read_write image1d_array_t scalarOCLImage1dArrayRW);
 void kernel OCLImage1dBufferRWTest(read_write image1d_buffer_t scalarOCLImage1dBufferRW);
 void kernel OCLImage2dRWTest(read_write image2d_t scalarOCLImage2dRW);
 void kernel OCLImage2dArrayRWTest(read_write image2d_array_t scalarOCLImage2dArrayRW);
 void kernel OCLImage2dDepthRWTest(read_write image2d_depth_t scalarOCLImage2dDepthRW);
 void kernel OCLImage2dArrayDepthRWTest(read_write image2d_array_depth_t scalarOCLImage2dArrayDepthRW);
 void kernel OCLImage2dMSAARWTest(read_write image2d_msaa_t scalarOCLImage2dMSAARW);
 void kernel OCLImage2dArrayMSAARWTest(read_write image2d_array_msaa_t scalarOCLImage2dArrayMSAARW);
 void kernel OCLImage2dMSAADepthRWTest(read_write image2d_msaa_depth_t scalarOCLImage2dMSAADepthRW);
 void kernel OCLImage2dArrayMSAADepthRWTest(read_write image2d_array_msaa_depth_t scalarOCLImage2dArrayMSAADepthRW);
 void kernel OCLImage3dRWTest(read_write image3d_t scalarOCLImage3dRW);
 
-// CHECK: ParmDecl=scalarOCLImage1dRW:80:51 (Definition) [type=__read_write image1d_t] [typekind=OCLImage1dRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage1dArrayRW:81:62 (Definition) [type=__read_write image1d_array_t] [typekind=OCLImage1dArrayRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage1dBufferRW:82:64 (Definition) [type=__read_write image1d_buffer_t] [typekind=OCLImage1dBufferRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dRW:83:51 (Definition) [type=__read_write image2d_t] [typekind=OCLImage2dRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayRW:84:62 (Definition) [type=__read_write image2d_array_t] [typekind=OCLImage2dArrayRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dDepthRW:85:62 (Definition) [type=__read_write image2d_depth_t] [typekind=OCLImage2dDepthRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayDepthRW:86:73 (Definition) [type=__read_write image2d_array_depth_t] [typekind=OCLImage2dArrayDepthRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dMSAARW:87:60 (Definition){{( \(invalid\))?}} [type=__read_write image2d_msaa_t] [typekind=OCLImage2dMSAARW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayMSAARW:88:71 (Definition){{( \(invalid\))?}} [type=__read_write image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAARW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dMSAADepthRW:89:71 (Definition){{( \(invalid\))?}} [type=__read_write image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthRW:90:82 (Definition){{( \(invalid\))?}} [type=__read_write image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthRW] [isPOD=1]
-// CHECK: ParmDecl=scalarOCLImage3dRW:91:51 (Definition) [type=__read_write image3d_t] [typekind=OCLImage3dRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dRW:80:51 (Definition) [type=__private __read_write image1d_t] [typekind=OCLImage1dRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dArrayRW:81:62 (Definition) [type=__private __read_write image1d_array_t] [typekind=OCLImage1dArrayRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage1dBufferRW:82:64 (Definition) [type=__private __read_write image1d_buffer_t] [typekind=OCLImage1dBufferRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dRW:83:51 (Definition) [type=__private __read_write image2d_t] [typekind=OCLImage2dRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayRW:84:62 (Definition) [type=__private __read_write image2d_array_t] [typekind=OCLImage2dArrayRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dDepthRW:85:62 (Definition) [type=__private __read_write image2d_depth_t] [typekind=OCLImage2dDepthRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayDepthRW:86:73 (Definition) [type=__private __read_write image2d_array_depth_t] [typekind=OCLImage2dArrayDepthRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dMSAARW:87:60 (Definition){{( \(invalid\))?}} [type=__private __read_write image2d_msaa_t] [typekind=OCLImage2dMSAARW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayMSAARW:88:71 (Definition){{( \(invalid\))?}} [type=__private __read_write image2d_array_msaa_t] [typekind=OCLImage2dArrayMSAARW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dMSAADepthRW:89:71 (Definition){{( \(invalid\))?}} [type=__private __read_write image2d_msaa_depth_t] [typekind=OCLImage2dMSAADepthRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage2dArrayMSAADepthRW:90:82 (Definition){{( \(invalid\))?}} [type=__private __read_write image2d_array_msaa_depth_t] [typekind=OCLImage2dArrayMSAADepthRW] [isPOD=1]
+// CHECK: ParmDecl=scalarOCLImage3dRW:91:51 (Definition) [type=__private __read_write image3d_t] [typekind=OCLImage3dRW] [isPOD=1]
 
 void kernel intPipeTestRO(read_only pipe int scalarPipe);
 void kernel intPipeTestWO(write_only pipe int scalarPipe);
 
-// CHECK: ParmDecl=scalarPipe:106:46 (Definition) [type=read_only pipe int] [typekind=Pipe] [isPOD=0]
-// CHECK: ParmDecl=scalarPipe:107:47 (Definition) [type=write_only pipe int] [typekind=Pipe] [isPOD=0]
+// CHECK: ParmDecl=scalarPipe:106:46 (Definition) [type=__private read_only pipe int] [typekind=Pipe] [isPOD=0]
+// CHECK: ParmDecl=scalarPipe:107:47 (Definition) [type=__private write_only pipe int] [typekind=Pipe] [isPOD=0]
 
 #define CLK_ADDRESS_CLAMP_TO_EDGE       2
 #define CLK_NORMALIZED_COORDS_TRUE      1
 #define CLK_FILTER_NEAREST              0x10
 
 void kernel testMiscOpenCLTypes() {
   const sampler_t scalarOCLSampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
   clk_event_t scalarOCLEvent;
   queue_t scalarOCLQueue;
   reserve_id_t scalarOCLReserveID;
 }
 
 // CHECK: VarDecl=scalarOCLSampler:117:19 (Definition) [type=const sampler_t] [typekind=Typedef] const [canonicaltype=const sampler_t] [canonicaltypekind=OCLSampler] [isPOD=1]
-// CHECK: VarDecl=scalarOCLEvent:118:15 (Definition) [type=clk_event_t] [typekind=Typedef] [canonicaltype=clk_event_t] [canonicaltypekind=Unexposed] [isPOD=1]
-// CHECK: VarDecl=scalarOCLQueue:119:11 (Definition) [type=queue_t] [typekind=Typedef] [canonicaltype=queue_t] [canonicaltypekind=OCLQueue] [isPOD=1]
-// CHECK: VarDecl=scalarOCLReserveID:120:16 (Definition) [type=reserve_id_t] [typekind=Typedef] [canonicaltype=reserve_id_t] [canonicaltypekind=OCLReserveID] [isPOD=1]
+// CHECK: VarDecl=scalarOCLEvent:118:15 (Definition) [type=__private clk_event_t] [typekind=Typedef] [canonicaltype=__private clk_event_t] [canonicaltypekind=Unexposed] [isPOD=1]
+// CHECK: VarDecl=scalarOCLQueue:119:11 (Definition) [type=__private queue_t] [typekind=Typedef] [canonicaltype=__private queue_t] [canonicaltypekind=OCLQueue] [isPOD=1]
+// CHECK: VarDecl=scalarOCLReserveID:120:16 (Definition) [type=__private reserve_id_t] [typekind=Typedef] [canonicaltype=__private reserve_id_t] [canonicaltypekind=OCLReserveID] [isPOD=1]
 
 #pragma OPENCL EXTENSION cl_intel_device_side_avc_motion_estimation : enable
 
 void kernel testExtOpenCLTypes() {
   intel_sub_group_avc_mce_payload_t mce_payload;
 }
 
-// CHECK: VarDecl=mce_payload:131:37 (Definition){{( \(invalid\))?}} [type=intel_sub_group_avc_mce_payload_t] [typekind=Typedef] [canonicaltype=intel_sub_group_avc_mce_payload_t] [canonicaltypekind=OCLIntelSubgroupAVCMcePayload] [isPOD=1]
+// CHECK: VarDecl=mce_payload:131:37 (Definition){{( \(invalid\))?}} [type=__private intel_sub_group_avc_mce_payload_t] [typekind=Typedef] [canonicaltype=__private intel_sub_group_avc_mce_payload_t] [canonicaltypekind=OCLIntelSubgroupAVCMcePayload] [isPOD=1]

clang/test/Parser/opencl-astype.cl

 // RUN: %clang_cc1 -fsyntax-only -verify %s -triple spir-unknown-unknown
 #pragma OPENCL EXTENSION cl_khr_fp64 : enable
 
 void test_astype() {
   float f = 1.0f;
   unsigned int i = __builtin_astype(f, unsigned int); 
   
   typedef __attribute__(( ext_vector_type(4) ))  int int4;
   typedef __attribute__(( ext_vector_type(3) ))  float float3;
   typedef __attribute__(( ext_vector_type(4) ))  float float4;
   typedef __attribute__(( ext_vector_type(4) ))  double double4;
   
   float4 f4;
-  double4 d4 = __builtin_astype(f4, double4); // expected-error{{invalid reinterpretation: sizes of 'double4' (vector of 4 'double' values) and 'float4' (vector of 4 'float' values) must match}}
+  double4 d4 = __builtin_astype(f4, double4); // expected-error{{invalid reinterpretation: sizes of 'double4' (vector of 4 'double' values) and '__private float4' (vector of 4 'float' values) must match}}
   
   // Verify int4->float3, float3->int4 works.
   int4 i4;
   float3 f3 = __builtin_astype(i4, float3);
   i4 = __builtin_astype(f3, int4);
 }

clang/test/Parser/opencl-atomics-cl20.cl

 #endif
 
 #ifdef CL20
 void foo(atomic_int * ptr) {}
 void atomic_ops_test() {
   atomic_int i;
   foo(&i);
 // OpenCL v2.0 s6.13.11.8, arithemtic operations are not permitted on atomic types.
-  i++; // expected-error {{invalid argument type 'atomic_int' (aka '_Atomic(int)') to unary expression}}
+  i++; // expected-error {{invalid argument type '__private atomic_int' (aka '__private _Atomic(int)') to unary expression}}
   i = 1; // expected-error {{atomic variable can be assigned to a variable only in global address space}}
-  i += 1; // expected-error {{invalid operands to binary expression ('atomic_int' (aka '_Atomic(int)') and 'int')}}
-  i = i + i; // expected-error {{invalid operands to binary expression ('atomic_int' (aka '_Atomic(int)') and 'atomic_int')}}
+  i += 1; // expected-error {{invalid operands to binary expression ('__private atomic_int' (aka '__private _Atomic(int)') and 'int')}}
+  i = i + i; // expected-error {{invalid operands to binary expression ('__private atomic_int' (aka '__private _Atomic(int)') and '__private atomic_int')}}
 }
 #endif

clang/test/SemaOpenCL/access-qualifier.cl

 #if __OPENCL_C_VERSION__ >= 200
 void myReadWrite(read_write image1d_t);
 #else
 void myReadWrite(read_write image1d_t); // expected-error {{access qualifier 'read_write' can not be used for '__read_write image1d_t' prior to OpenCL version 2.0}}
 #endif
 
 
 kernel void k1(img1d_wo img) {
-  myRead(img); // expected-error {{passing 'img1d_wo' (aka '__write_only image1d_t') to parameter of incompatible type '__read_only image1d_t'}}
+  myRead(img); // expected-error {{passing '__private img1d_wo' (aka '__private __write_only image1d_t') to parameter of incompatible type '__read_only image1d_t'}}
 }
 
 kernel void k2(img1d_ro img) {
-  myWrite(img); // expected-error {{passing 'img1d_ro' (aka '__read_only image1d_t') to parameter of incompatible type '__write_only image1d_t'}}
+  myWrite(img); // expected-error {{passing '__private img1d_ro' (aka '__private __read_only image1d_t') to parameter of incompatible type '__write_only image1d_t'}}
 }
 
 kernel void k3(img1d_wo img) {
   myWrite(img);
 }
 
 #if __OPENCL_C_VERSION__ >= 200
 kernel void k4(img1d_rw img) {
   myReadWrite(img);
 }
 #endif
 
 kernel void k5(img1d_ro_default img) {
-  myWrite(img); // expected-error {{passing 'img1d_ro_default' (aka '__read_only image1d_t') to parameter of incompatible type '__write_only image1d_t'}}
+  myWrite(img); // expected-error {{passing '__private img1d_ro_default' (aka '__private __read_only image1d_t') to parameter of incompatible type '__write_only image1d_t'}}
 }
 
 kernel void k6(img1d_ro img) {
   myRead(img);
 }
 
 kernel void k7(read_only img1d_wo img){} // expected-error {{multiple access qualifiers}}
 

 kernel void k13(read_write pipe int i){} // expected-error{{access qualifier 'read_write' can not be used for 'read_only pipe int'}}
 #else
 kernel void k13(__read_write image1d_t i){} // expected-error{{access qualifier '__read_write' can not be used for '__read_write image1d_t' prior to OpenCL version 2.0}}
 #endif
 
 #if __OPENCL_C_VERSION__ >= 200
 void myPipeWrite(write_only pipe int); // expected-note {{passing argument to parameter here}}
 kernel void k14(read_only pipe int p) {
-  myPipeWrite(p); // expected-error {{passing 'read_only pipe int' to parameter of incompatible type 'write_only pipe int'}}
+  myPipeWrite(p); // expected-error {{passing '__private read_only pipe int' to parameter of incompatible type 'write_only pipe int'}}
 }
 #endif
 
 #if __OPENCL_C_VERSION__ < 200
 kernel void test_image3d_wo(write_only image3d_t img) {} // expected-error {{use of type '__write_only image3d_t' requires cl_khr_3d_image_writes extension to be enabled}}
 #endif
 
 #if __OPENCL_C_VERSION__ >= 200
 kernel void read_write_twice_typedef(read_write img1d_rw i){} // expected-warning {{duplicate 'read_write' declaration specifier}}
 // expected-note@-74 {{previously declared 'read_write' here}}
 
 kernel void pipe_ro_twice(read_only read_only pipe int i){} // expected-warning{{duplicate 'read_only' declaration specifier}}
 // Conflicting access qualifiers
 kernel void pipe_ro_twice_tw(read_write read_only read_only pipe int i){} // expected-error{{access qualifier 'read_write' can not be used for 'read_only pipe int'}}
 kernel void pipe_ro_wo(read_only write_only pipe int i){} // expected-error{{multiple access qualifiers}}
 
 typedef read_only pipe int ROPipeInt;
 kernel void pipe_ro_twice_typedef(read_only ROPipeInt i){} // expected-warning{{duplicate 'read_only' declaration specifier}}
 // expected-note@-2 {{previously declared 'read_only' here}}
 
 kernel void pass_ro_typedef_to_wo(ROPipeInt p) {
-  myPipeWrite(p); // expected-error {{passing 'ROPipeInt' (aka 'read_only pipe int') to parameter of incompatible type 'write_only pipe int'}}
+  myPipeWrite(p); // expected-error {{passing '__private ROPipeInt' (aka '__private read_only pipe int') to parameter of incompatible type 'write_only pipe int'}}
   // expected-note@-25 {{passing argument to parameter here}}
 }
 #endif
 
 kernel void read_only_twice_typedef(__read_only img1d_ro i){} // expected-warning {{duplicate '__read_only' declaration specifier}}
 // expected-note@-95 {{previously declared 'read_only' here}}
 
 kernel void read_only_twice_default(read_only img1d_ro_default img){} // expected-warning {{duplicate 'read_only' declaration specifier}}
 // expected-note@-101 {{previously declared 'read_only' here}}
 
 kernel void image_wo_twice(write_only __write_only image1d_t i){} // expected-warning {{duplicate '__write_only' declaration specifier}}
 kernel void image_wo_twice_typedef(write_only img1d_wo i){} // expected-warning {{duplicate 'write_only' declaration specifier}}
 // expected-note@-103 {{previously declared 'write_only' here}}
 

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

 // expected-note-re@-5{{candidate function not viable: cannot pass pointer to address space '__{{global|generic}}' as a pointer to address space '__constant' in 1st argument}}
 #endif
 #endif
 
 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: cannot pass pointer to address space '__{{global|generic|constant}}' as a pointer to default address space in 1st argument}}
+// expected-note-re@-4{{candidate function not viable: cannot pass pointer to address space '__{{global|generic|constant}}' as a pointer to address space '__private' in 1st argument}}
 #endif
 
 void f_gen(__generic int *arg_gen) {}
 #ifdef CONSTANT
 #if !__OPENCL_CPP_VERSION__
 // expected-note@-3{{passing argument to parameter 'arg_gen' here}}
 #else
 // expected-note@-5{{candidate function not viable: cannot pass pointer to address space '__constant' as a pointer to address space '__generic' in 1st argument}}
 #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) {
 
   AS int *var_init1 = arg_glob;
 #ifdef CONSTANT
 #if !__OPENCL_CPP_VERSION__
-// expected-error@-3{{initializing '__constant int *' with an expression of type '__global int *' changes address space of pointer}}
+// expected-error@-3{{initializing '__constant int *__private' with an expression of type '__global int *__private' changes address space of pointer}}
 #else
-// expected-error@-5{{cannot initialize a variable of type '__constant int *' with an lvalue of type '__global int *'}}
+// expected-error@-5{{cannot initialize a variable of type '__constant int *__private' with an lvalue of type '__global int *__private'}}
 #endif
 #endif
 
   AS int *var_init2 = arg_loc;
 #ifndef GENERIC
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re@-3{{initializing '__{{global|constant}} int *' with an expression of type '__local int *' changes address space of pointer}}
+// expected-error-re@-3{{initializing '__{{global|constant}} int *__private' with an expression of type '__local int *__private' 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 *'}}
+// expected-error-re@-5{{cannot initialize a variable of type '__{{global|constant}} int *__private' with an lvalue of type '__local int *__private'}}
 #endif
 #endif
 
   AS int *var_init3 = arg_const;
 #ifndef CONSTANT
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re@-3{{initializing '__{{global|generic}} int *' with an expression of type '__constant int *' changes address space of pointer}}
+// expected-error-re@-3{{initializing '__{{global|generic}} int *__private' with an expression of type '__constant int *__private' 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 *'}}
+// expected-error-re@-5{{cannot initialize a variable of type '__{{global|generic}} int *__private' with an lvalue of type '__constant int *__private'}}
 #endif
 #endif
 
   AS int *var_init4 = arg_priv;
 #ifndef GENERIC
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re@-3{{initializing '__{{global|constant}} int *' with an expression of type 'int *' changes address space of pointer}}
+// expected-error-re@-3{{initializing '__{{global|constant}} int *__private' with an expression of type '__private int *__private' 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 *'}}
+// expected-error-re@-5{{cannot initialize a variable of type '__{{global|constant}} int *__private' with an lvalue of type '__private int *__private'}}
 #endif
 #endif
 
   AS int *var_init5 = arg_gen;
 #ifndef GENERIC
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re@-3{{initializing '__{{global|constant}} int *' with an expression of type '__generic int *' changes address space of pointer}}
+// expected-error-re@-3{{initializing '__{{global|constant}} int *__private' with an expression of type '__generic int *__private' 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 *'}}
+// expected-error-re@-5{{cannot initialize a variable of type '__{{global|constant}} int *__private' with an lvalue of type '__generic int *__private'}}
 #endif
 #endif
 
   AS int *var_cast1 = (AS int *)arg_glob;
 #ifdef CONSTANT
 #if !__OPENCL_CPP_VERSION__
 // expected-error@-3{{casting '__global int *' to type '__constant int *' changes address space of pointer}}
 #else

 #else
 // expected-error-re@-5{{C-style cast from '__constant int *' to '__{{global|generic}} int *' converts between mismatching address spaces}}
 #endif
 #endif
 
   AS int *var_cast4 = (AS int *)arg_priv;
 #ifndef GENERIC
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re@-3{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-3{{casting '__private int *' to type '__{{global|constant}} int *' changes address space of pointer}}
 #else
-// expected-error-re@-5{{C-style cast from 'int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
+// expected-error-re@-5{{C-style cast from '__private int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
 #endif
 #endif
 
   AS int *var_cast5 = (AS int *)arg_gen;
 #ifdef CONSTANT
 #if !__OPENCL_CPP_VERSION__
 // expected-error@-3{{casting '__generic int *' to type '__constant int *' changes address space of pointer}}
 #else
 // expected-error@-5{{C-style cast from '__generic int *' to '__constant int *' converts between mismatching address spaces}}
 #endif
 #endif
 
   AS int *var_impl;
   var_impl = arg_glob;
 #ifdef CONSTANT
 #if !__OPENCL_CPP_VERSION__
-// expected-error@-3{{assigning '__global int *' to '__constant int *' changes address space of pointer}}
+// expected-error@-3{{assigning '__global int *__private' to '__constant int *__private' changes address space of pointer}}
 #else
-// expected-error@-5{{assigning to '__constant int *' from incompatible type '__global int *'}}
+// expected-error@-5{{assigning to '__constant int *' from incompatible type '__global int *__private'}}
 #endif
 #endif
 
   var_impl = arg_loc;
 #ifndef GENERIC
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re@-3{{assigning '__local int *' to '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-3{{assigning '__local int *__private' to '__{{global|constant}} int *__private' changes address space of pointer}}
 #else
-// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type '__local int *'}}
+// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type '__local int *__private'}}
 #endif
 #endif
 
   var_impl = arg_const;
 #ifndef CONSTANT
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re@-3{{assigning '__constant int *' to '__{{global|generic}} int *' changes address space of pointer}}
+// expected-error-re@-3{{assigning '__constant int *__private' to '__{{global|generic}} int *__private' changes address space of pointer}}
 #else
-// expected-error-re@-5{{assigning to '__{{global|generic}} int *' from incompatible type '__constant int *'}}
+// expected-error-re@-5{{assigning to '__{{global|generic}} int *' from incompatible type '__constant int *__private'}}
 #endif
 #endif
 
   var_impl = arg_priv;
 #ifndef GENERIC
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re@-3{{assigning 'int *' to '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-3{{assigning '__private int *__private' to '__{{global|constant}} int *__private' changes address space of pointer}}
 #else
-// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type 'int *'}}
+// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type '__private int *__private'}}
 #endif
 #endif
 
   var_impl = arg_gen;
 #ifndef GENERIC
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re@-3{{assigning '__generic int *' to '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-3{{assigning '__generic int *__private' to '__{{global|constant}} int *__private' changes address space of pointer}}
 #else
-// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type '__generic int *'}}
+// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type '__generic int *__private'}}
 #endif
 #endif
 
   var_cast1 = (AS int *)arg_glob;
 #ifdef CONSTANT
 #if !__OPENCL_CPP_VERSION__
 // expected-error@-3{{casting '__global int *' to type '__constant int *' changes address space of pointer}}
 #else

 #else
 // expected-error-re@-5{{C-style cast from '__constant int *' to '__{{global|generic}} int *' converts between mismatching address spaces}}
 #endif
 #endif
 
   var_cast4 = (AS int *)arg_priv;
 #ifndef GENERIC
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re@-3{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-3{{casting '__private int *' to type '__{{global|constant}} int *' changes address space of pointer}}
 #else
-// expected-error-re@-5{{C-style cast from 'int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
+// expected-error-re@-5{{C-style cast from '__private int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
 #endif
 #endif
 
   var_cast5 = (AS int *)arg_gen;
 #ifdef CONSTANT
 #if !__OPENCL_CPP_VERSION__
 // expected-error@-3{{casting '__generic int *' to type '__constant int *' changes address space of pointer}}
 #else

 #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
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re@-3{{comparison between  ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-3{{comparison between  ('__{{global|constant}} int *' and '__private int *') which are pointers to non-overlapping address spaces}}
 #else
-// expected-error-re@-5{{comparison of distinct pointer types ('__{{global|constant}} int *' and 'int *')}}
+// expected-error-re@-5{{comparison of distinct pointer types ('__{{global|constant}} int *' and '__private int *')}}
 #endif
 #endif
 
   b = var_cmp >= arg_gen;
 #ifdef CONSTANT
 #if !__OPENCL_CPP_VERSION__
 // expected-error@-3{{comparison between  ('__constant int *' and '__generic int *') which are pointers to non-overlapping address spaces}}
 #else
#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
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re@-3{{passing '__{{constant|generic}} int *' to parameter of type '__global int *' changes address space of pointer}}
+// expected-error-re@-3{{passing '__{{constant|generic}} int *__private' 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);
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re@-2{{passing '__{{global|constant|generic}} int *' to parameter of type '__local int *' changes address space of pointer}}
+// expected-error-re@-2{{passing '__{{global|constant|generic}} int *__private' 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
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re@-3{{passing '__{{global|generic}} int *' to parameter of type '__constant int *' changes address space of pointer}}
+// expected-error-re@-3{{passing '__{{global|generic}} int *__private' 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);
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re@-2{{passing '__{{global|constant|generic}} int *' to parameter of type 'int *' changes address space of pointer}}
+// expected-error-re@-2{{passing '__{{global|constant|generic}} int *__private' to parameter of type '__private 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
 #if !__OPENCL_CPP_VERSION__
-// expected-error@-3{{passing '__constant int *' to parameter of type '__generic int *' changes address space of pointer}}
+// expected-error@-3{{passing '__constant int *__private' 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;

 // expected-error-re@-5{{incompatible operand types ('__{{global|generic}} int *' and '__constant int *')}}
 #endif
 #endif
 
   __private int *var_priv;
   var_gen = 0 ? var_cond : var_priv;
 #ifndef GENERIC
 #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}}
+// expected-error-re@-3{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and '__private int *') which are pointers to non-overlapping address spaces}}
 #else
-// expected-error-re@-5{{incompatible operand types ('__{{global|constant}} int *' and 'int *')}}
+// expected-error-re@-5{{incompatible operand types ('__{{global|constant}} int *' and '__private int *')}}
 #endif
 #endif
 
   var_gen = 0 ? var_cond : var_gen;
 #ifdef CONSTANT
 #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

 #else
 // expected-warning@-7{{pointer type mismatch ('__constant int *' and '__constant char *')}}
 #endif
 #endif
 
   __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 *')}}
+// expected-error-re@-2{{incompatible operand types ('__{{constant|global|generic}} int *' and '__private char *')}}
 #else
 #ifndef GENERIC
-// 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}}
+// expected-error-re@-5{{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@-7{{pointer type mismatch ('__generic int *' and 'char *')}}
+// expected-warning@-7{{pointer type mismatch ('__generic int *' and '__private char *')}}
 #endif
 #endif
 
   __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

 // 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.
   // FIXME: Should this really be allowed in C++ mode?
   var_as_as_int = var_asc_asc_int;
 #if !__OPENCL_CPP_VERSION__
 #ifdef GENERIC
-// expected-error@-3 {{assigning '__local int *__local *' to '__generic int *__generic *' changes address space of nested pointer}}
+// expected-error@-3 {{assigning '__local int *__local *__private' to '__generic int *__generic *__private' changes address space of nested pointer}}
 #endif
 #endif
   var_as_as_int = 0 ? var_as_as_int : var_asc_asc_int;
 #if !__OPENCL_CPP_VERSION__
 #ifdef GENERIC
 // expected-warning@-3{{pointer type mismatch ('__generic int *__generic *' and '__local int *__local *')}}
 #endif
 #endif
 }

clang/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=clc++ -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;
 #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}}
+  ip = gip; // expected-error {{assigning '__global int *__private' to '__private int *__private' changes address space of pointer}}
+  ip = &li; // expected-error {{assigning '__local int *' to '__private int *__private' changes address space of pointer}}
+  ip = &ci; // expected-error {{assigning '__constant int *' to '__private int *__private' changes address space of pointer}}
 #else
   ip = gip;
   ip = &li;
   ip = &ci;
 #if !__OPENCL_CPP_VERSION__
-// expected-error@-2 {{assigning '__constant int *' to '__generic int *' changes address space of pointer}}
+// expected-error@-2 {{assigning '__constant int *' to '__generic int *__private' 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;
#endif
   g = (__global int *)cc;
 #if !__OPENCL_CPP_VERSION__
 // expected-error@-2 {{casting 'const __constant int *' to type '__global int *' changes address space of pointer}}
 #else
 // expected-error@-4 {{C-style cast from 'const __constant int *' to '__global int *' converts between mismatching address spaces}}
 #endif
   g = (__global int *)p;
 #if !__OPENCL_CPP_VERSION__
-// expected-error@-2 {{casting 'int *' to type '__global int *' changes address space of pointer}}
+// expected-error@-2 {{casting '__private int *' to type '__global int *' changes address space of pointer}}
 #else
-// expected-error@-4 {{C-style cast from 'int *' to '__global int *' converts between mismatching address spaces}}
+// expected-error@-4 {{C-style cast from '__private int *' to '__global int *' converts between mismatching address spaces}}
 #endif
   l = (__local int *)g;
 #if !__OPENCL_CPP_VERSION__
 // expected-error@-2 {{casting '__global int *' to type '__local int *' changes address space of pointer}}
 #else
 // expected-error@-4 {{C-style cast from '__global int *' to '__local int *' converts between mismatching address spaces}}
 #endif
   l = (__local int *)c;
 #if !__OPENCL_CPP_VERSION__
 // expected-error@-2 {{casting '__constant int *' to type '__local int *' changes address space of pointer}}
 #else
 // expected-error@-4 {{C-style cast from '__constant int *' to '__local int *' converts between mismatching address spaces}}
 #endif
   l = (__local int *)cc;
 #if !__OPENCL_CPP_VERSION__
 // expected-error@-2 {{casting 'const __constant int *' to type '__local int *' changes address space of pointer}}
 #else
 // expected-error@-4 {{C-style cast from 'const __constant int *' to '__local int *' converts between mismatching address spaces}}
 #endif
   l = (__local int *)p;
 #if !__OPENCL_CPP_VERSION__
-// expected-error@-2 {{casting 'int *' to type '__local int *' changes address space of pointer}}
+// expected-error@-2 {{casting '__private int *' to type '__local int *' changes address space of pointer}}
 #else
-// expected-error@-4 {{C-style cast from 'int *' to '__local int *' converts between mismatching address spaces}}
+// expected-error@-4 {{C-style cast from '__private int *' to '__local int *' converts between mismatching address spaces}}
 #endif
   c = (__constant int *)g;
 #if !__OPENCL_CPP_VERSION__
 // expected-error@-2 {{casting '__global int *' to type '__constant int *' changes address space of pointer}}
 #else
 // expected-error@-4 {{C-style cast from '__global int *' to '__constant int *' converts between mismatching address spaces}}
 #endif
   c = (__constant int *)l;
 #if !__OPENCL_CPP_VERSION__
 // expected-error@-2 {{casting '__local int *' to type '__constant int *' changes address space of pointer}}
 #else
 // expected-error@-4 {{C-style cast from '__local int *' to '__constant int *' converts between mismatching address spaces}}
 #endif
   c = (__constant int *)p;
 #if !__OPENCL_CPP_VERSION__
-// expected-error@-2 {{casting 'int *' to type '__constant int *' changes address space of pointer}}
+// expected-error@-2 {{casting '__private int *' to type '__constant int *' changes address space of pointer}}
 #else
-// expected-error@-4 {{C-style cast from 'int *' to '__constant int *' converts between mismatching address spaces}}
+// expected-error@-4 {{C-style cast from '__private int *' to '__constant int *' converts between mismatching address spaces}}
 #endif
   p = (__private int *)g;
 #if !__OPENCL_CPP_VERSION__
-// expected-error@-2 {{casting '__global int *' to type 'int *' changes address space of pointer}}
+// expected-error@-2 {{casting '__global int *' to type '__private int *' changes address space of pointer}}
 #else
-// expected-error@-4 {{C-style cast from '__global int *' to 'int *' converts between mismatching address spaces}}
+// expected-error@-4 {{C-style cast from '__global int *' to '__private int *' converts between mismatching address spaces}}
 #endif
   p = (__private int *)l;
 #if !__OPENCL_CPP_VERSION__
-// expected-error@-2 {{casting '__local int *' to type 'int *' changes address space of pointer}}
+// expected-error@-2 {{casting '__local int *' to type '__private int *' changes address space of pointer}}
 #else
-// expected-error@-4 {{C-style cast from '__local int *' to 'int *' converts between mismatching address spaces}}
+// expected-error@-4 {{C-style cast from '__local int *' to '__private int *' converts between mismatching address spaces}}
 #endif
   p = (__private int *)c;
 #if !__OPENCL_CPP_VERSION__
-// expected-error@-2 {{casting '__constant int *' to type 'int *' changes address space of pointer}}
+// expected-error@-2 {{casting '__constant int *' to type '__private int *' changes address space of pointer}}
 #else
-// expected-error@-4 {{C-style cast from '__constant int *' to 'int *' converts between mismatching address spaces}}
+// expected-error@-4 {{C-style cast from '__constant int *' to '__private int *' converts between mismatching address spaces}}
 #endif
   p = (__private int *)cc;
 #if !__OPENCL_CPP_VERSION__
-// expected-error@-2 {{casting 'const __constant int *' to type 'int *' changes address space of pointer}}
+// expected-error@-2 {{casting 'const __constant int *' to type '__private int *' changes address space of pointer}}
 #else
-// expected-error@-4 {{C-style cast from 'const __constant int *' to 'int *' converts between mismatching address spaces}}
+// expected-error@-4 {{C-style cast from 'const __constant int *' to '__private int *' converts between mismatching address spaces}}
 #endif
 }
 
 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;
 }
 
 #if !__OPENCL_CPP_VERSION__
 void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) {
-  g = gg;    // expected-error {{assigning '__global int **' to '__global int *' changes address space of pointer}}
-  g = l;     // expected-error {{assigning '__local int *' to '__global int *' changes address space of pointer}}
-  g = ll;    // expected-error {{assigning '__local int **' to '__global int *' changes address space of pointer}}
-  g = gg_f;  // expected-error {{assigning '__global float **' to '__global int *' changes address space of pointer}}
-  g = (__global int *)gg_f; // expected-error {{casting '__global float **' to type '__global int *' changes address space of pointer}}
-
-  gg = g;    // expected-error {{assigning '__global int *' to '__global int **' changes address space of pointer}}
-  gg = l;    // expected-error {{assigning '__local int *' to '__global int **' changes address space of pointer}}
-  gg = ll;   // expected-error {{assigning '__local int **' to '__global int **' changes address space of nested pointer}}
-  gg = gg_f; // expected-warning {{incompatible pointer types assigning to '__global int **' from '__global float **'}}
+  g = gg;    // expected-error {{assigning '__global int *__private *__private' to '__global int *__private' changes address space of pointer}}
+  g = l;     // expected-error {{assigning '__local int *__private' to '__global int *__private' changes address space of pointer}}
+  g = ll;    // expected-error {{assigning '__local int *__private *__private' to '__global int *__private' changes address space of pointer}}
+  g = gg_f;  // expected-error {{assigning '__global float *__private *__private' to '__global int *__private' changes address space of pointer}}
+  g = (__global int *)gg_f; // expected-error {{casting '__global float *__private *' to type '__global int *' changes address space of pointer}}
+
+  gg = g;    // expected-error {{assigning '__global int *__private' to '__global int *__private *__private' changes address space of pointer}}
+  gg = l;    // expected-error {{assigning '__local int *__private' to '__global int *__private *__private' changes address space of pointer}}
+  gg = ll;   // expected-error {{assigning '__local int *__private *__private' to '__global int *__private *__private' changes address space of nested pointer}}
+  gg = gg_f; // expected-warning {{incompatible pointer types assigning to '__global int *__private *__private' from '__global float *__private *__private'}}
   gg = (__global int * __private *)gg_f;
 
-  l = g;     // expected-error {{assigning '__global int *' to '__local int *' changes address space of pointer}}
-  l = gg;    // expected-error {{assigning '__global int **' to '__local int *' changes address space of pointer}}
-  l = ll;    // expected-error {{assigning '__local int **' to '__local int *' changes address space of pointer}}
-  l = gg_f;  // expected-error {{assigning '__global float **' to '__local int *' changes address space of pointer}}
-  l = (__local int *)gg_f; // expected-error {{casting '__global float **' to type '__local int *' changes address space of pointer}}
-
-  ll = g;    // expected-error {{assigning '__global int *' to '__local int **' changes address space of pointer}}
-  ll = gg;   // expected-error {{assigning '__global int **' to '__local int **' changes address space of nested pointer}}
-  ll = l;    // expected-error {{assigning '__local int *' to '__local int **' changes address space of pointer}}
-  ll = gg_f; // expected-error {{assigning '__global float **' to '__local int **' changes address space of nested pointer}}
-  ll = (__local int * __private *)gg_f; // expected-warning {{casting '__global float **' to type '__local int **' discards qualifiers in nested pointer types}}
-
-  gg_f = g;  // expected-error {{assigning '__global int *' to '__global float **' changes address space of pointer}}
-  gg_f = gg; // expected-warning {{incompatible pointer types assigning to '__global float **' from '__global int **'}}
-  gg_f = l;  // expected-error {{assigning '__local int *' to '__global float **' changes address space of pointer}}
-  gg_f = ll; // expected-error {{assigning '__local int **' to '__global float **' changes address space of nested pointer}}
+  l = g;     // expected-error {{assigning '__global int *__private' to '__local int *__private' changes address space of pointer}}
+  l = gg;    // expected-error {{assigning '__global int *__private *__private' to '__local int *__private' changes address space of pointer}}
+  l = ll;    // expected-error {{assigning '__local int *__private *__private' to '__local int *__private' changes address space of pointer}}
+  l = gg_f;  // expected-error {{assigning '__global float *__private *__private' to '__local int *__private' changes address space of pointer}}
+  l = (__local int *)gg_f; // expected-error {{casting '__global float *__private *' to type '__local int *' changes address space of pointer}}
+
+  ll = g;    // expected-error {{assigning '__global int *__private' to '__local int *__private *__private' changes address space of pointer}}
+  ll = gg;   // expected-error {{assigning '__global int *__private *__private' to '__local int *__private *__private' changes address space of nested pointer}}
+  ll = l;    // expected-error {{assigning '__local int *__private' to '__local int *__private *__private' changes address space of pointer}}
+  ll = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__local int *__private *__private' changes address space of nested pointer}}
+  ll = (__local int * __private *)gg_f; // expected-warning {{casting '__global float *__private *' to type '__local int *__private *' discards qualifiers in nested pointer types}}
+
+  gg_f = g;  // expected-error {{assigning '__global int *__private' to '__global float *__private *__private' changes address space of pointer}}
+  gg_f = gg; // expected-warning {{incompatible pointer types assigning to '__global float *__private *__private' from '__global int *__private *__private'}}
+  gg_f = l;  // expected-error {{assigning '__local int *__private' to '__global float *__private *__private' changes address space of pointer}}
+  gg_f = ll; // expected-error {{assigning '__local int *__private *__private' to '__global float *__private *__private' changes address space of nested pointer}}
   gg_f = (__global float * __private *)gg;
 
   // FIXME: This doesn't seem right. This should be an error, not a warning.
   __local int * __global * __private * lll;
-  lll = gg; // expected-warning {{incompatible pointer types assigning to '__local int *__global **' from '__global int **'}}
+  lll = gg; // expected-warning {{incompatible pointer types assigning to '__local int *__global *__private *__private' from '__global int *__private *__private'}}
 
   typedef __local int * l_t;
   typedef __global int * g_t;
   __private l_t * pl;
   __private g_t * pg;
-  gg = pl;  // expected-error {{assigning 'l_t *' (aka '__local int **') to '__global int **' changes address space of nested pointer}}
-  pl = gg;  // expected-error {{assigning '__global int **' to 'l_t *' (aka '__local int **') changes address space of nested pointer}}
+  gg = pl;  // expected-error {{assigning '__private l_t *__private' (aka '__local int *__private *__private') to '__global int *__private *__private' changes address space of nested pointer}}
+  pl = gg;  // expected-error {{assigning '__global int *__private *__private' to '__private l_t *__private' (aka '__local int *__private *__private') changes address space of nested pointer}}
   gg = pg;
   pg = gg;
-  pg = pl;  // expected-error {{assigning 'l_t *' (aka '__local int **') to 'g_t *' (aka '__global int **') changes address space of nested pointer}}
-  pl = pg;  // expected-error {{assigning 'g_t *' (aka '__global int **') to 'l_t *' (aka '__local int **') changes address space of nested pointer}}
+  pg = pl;  // expected-error {{assigning '__private l_t *__private' (aka '__local int *__private *__private') to '__private g_t *__private' (aka '__global int *__private *__private') changes address space of nested pointer}}
+  pl = pg;  // expected-error {{assigning '__private g_t *__private' (aka '__global int *__private *__private') to '__private l_t *__private' (aka '__local int *__private *__private') changes address space of nested pointer}}
 
   ll = (__local int * __private *)(void *)gg;
   void *vp = ll;
 }
 #else
 void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) {
-  g = gg;    // expected-error {{assigning to '__global int *' from incompatible type '__global int **'}}
-  g = l;     // expected-error {{assigning to '__global int *' from incompatible type '__local int *'}}
-  g = ll;    // expected-error {{assigning to '__global int *' from incompatible type '__local int **'}}
-  g = gg_f;  // expected-error {{assigning to '__global int *' from incompatible type '__global float **'}}
-  g = (__global int *)gg_f; // expected-error {{C-style cast from '__global float **' to '__global int *' converts between mismatching address spaces}}
-
-  gg = g;    // expected-error {{assigning to '__global int **' from incompatible type '__global int *'; take the address with &}}
-  gg = l;    // expected-error {{assigning to '__global int **' from incompatible type '__local int *'}}
-  gg = ll;   // expected-error {{assigning to '__global int **' from incompatible type '__local int **'}}
-  gg = gg_f; // expected-error {{assigning to '__global int **' from incompatible type '__global float **'}}
+  g = gg;    // expected-error {{assigning to '__global int *' from incompatible type '__global int *__private *__private'}}
+  g = l;     // expected-error {{assigning to '__global int *' from incompatible type '__local int *__private'}}
+  g = ll;    // expected-error {{assigning to '__global int *' from incompatible type '__local int *__private *__private'}}
+  g = gg_f;  // expected-error {{assigning to '__global int *' from incompatible type '__global float *__private *__private'}}
+  g = (__global int *)gg_f; // expected-error {{C-style cast from '__global float *__private *' to '__global int *' converts between mismatching address spaces}}
+
+  gg = g;    // expected-error {{assigning to '__global int *__private *' from incompatible type '__global int *__private'; take the address with &}}
+  gg = l;    // expected-error {{assigning to '__global int *__private *' from incompatible type '__local int *__private'}}
+  gg = ll;   // expected-error {{assigning to '__global int *__private *' from incompatible type '__local int *__private *__private'}}
+  gg = gg_f; // expected-error {{assigning to '__global int *__private *' from incompatible type '__global float *__private *__private'}}
   gg = (__global int * __private *)gg_f;
 
-  l = g;     // expected-error {{assigning to '__local int *' from incompatible type '__global int *'}}
-  l = gg;    // expected-error {{assigning to '__local int *' from incompatible type '__global int **'}}
-  l = ll;    // expected-error {{assigning to '__local int *' from incompatible type '__local int **'}}
-  l = gg_f;  // expected-error {{assigning to '__local int *' from incompatible type '__global float **'}}
-  l = (__local int *)gg_f; // expected-error {{C-style cast from '__global float **' to '__local int *' converts between mismatching address spaces}}
-
-  ll = g;    // expected-error {{assigning to '__local int **' from incompatible type '__global int *'}}
-  ll = gg;   // expected-error {{assigning to '__local int **' from incompatible type '__global int **'}}
-  ll = l;    // expected-error {{assigning to '__local int **' from incompatible type '__local int *'; take the address with &}}
-  ll = gg_f; // expected-error {{assigning to '__local int **' from incompatible type '__global float **'}}
+  l = g;     // expected-error {{assigning to '__local int *' from incompatible type '__global int *__private'}}
+  l = gg;    // expected-error {{assigning to '__local int *' from incompatible type '__global int *__private *__private'}}
+  l = ll;    // expected-error {{assigning to '__local int *' from incompatible type '__local int *__private *__private'}}
+  l = gg_f;  // expected-error {{assigning to '__local int *' from incompatible type '__global float *__private *__private'}}
+  l = (__local int *)gg_f; // expected-error {{C-style cast from '__global float *__private *' to '__local int *' converts between mismatching address spaces}}
+
+  ll = g;    // expected-error {{assigning to '__local int *__private *' from incompatible type '__global int *__private'}}
+  ll = gg;   // expected-error {{assigning to '__local int *__private *' from incompatible type '__global int *__private *__private'}}
+  ll = l;    // expected-error {{assigning to '__local int *__private *' from incompatible type '__local int *__private'; take the address with &}}
+  ll = gg_f; // expected-error {{assigning to '__local int *__private *' from incompatible type '__global float *__private *__private'}}
   // FIXME: The below becomes a reinterpret_cast, and therefore does not emit an error
   // even though the address space mismatches in the nested pointers.
   ll = (__local int * __private *)gg;
 
-  gg_f = g;  // expected-error {{assigning to '__global float **' from incompatible type '__global int *'}}
-  gg_f = gg; // expected-error {{assigning to '__global float **' from incompatible type '__global int **'}}
-  gg_f = l;  // expected-error {{assigning to '__global float **' from incompatible type '__local int *'}}
-  gg_f = ll; // expected-error {{assigning to '__global float **' from incompatible type '__local int **'}}
+  gg_f = g;  // expected-error {{assigning to '__global float *__private *' from incompatible type '__global int *__private'}}
+  gg_f = gg; // expected-error {{assigning to '__global float *__private *' from incompatible type '__global int *__private *__private'}}
+  gg_f = l;  // expected-error {{assigning to '__global float *__private *' from incompatible type '__local int *__private'}}
+  gg_f = ll; // expected-error {{assigning to '__global float *__private *' from incompatible type '__local int *__private *__private'}}
   gg_f = (__global float * __private *)gg;
 
   typedef __local int * l_t;
   typedef __global int * g_t;
   __private l_t * pl;
   __private g_t * pg;
-  gg = pl;  // expected-error {{assigning to '__global int **' from incompatible type 'l_t *' (aka '__local int **')}}
-  pl = gg;  // expected-error {{assigning to 'l_t *' (aka '__local int **') from incompatible type '__global int **'}}
+  gg = pl;  // expected-error {{assigning to '__global int *__private *' from incompatible type '__private l_t *__private' (aka '__local int *__private *__private')}}
+  pl = gg;  // expected-error {{assigning to '__private l_t *' (aka '__local int *__private *') from incompatible type '__global int *__private *__private'}}
   gg = pg;
   pg = gg;
-  pg = pl;  // expected-error {{assigning to 'g_t *' (aka '__global int **') from incompatible type 'l_t *' (aka '__local int **')}}
-  pl = pg;  // expected-error {{assigning to 'l_t *' (aka '__local int **') from incompatible type 'g_t *' (aka '__global int **')}}
+  pg = pl;  // expected-error {{assigning to '__private g_t *' (aka '__global int *__private *') from incompatible type '__private l_t *__private' (aka '__local int *__private *__private')}}
+  pl = pg;  // expected-error {{assigning to '__private l_t *' (aka '__local int *__private *') from incompatible type '__private g_t *__private' (aka '__global int *__private *__private')}}
 
   ll = (__local int * __private *)(void *)gg;
   void *vp = ll;
 }
 #endif
 
 __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}}

clang/test/SemaOpenCL/arithmetic-conversions.cl

 // RUN: %clang_cc1 %s -triple spir-unknown-unknown -verify -pedantic -fsyntax-only -cl-std=CL1.2
 
 typedef float float2 __attribute__((ext_vector_type(2)));
 typedef long long2 __attribute__((ext_vector_type(2)));
 typedef int int2 __attribute__((ext_vector_type(2)));
 
 kernel void foo1(float2 in, global float2 *out) { *out = in + 0.5;} // expected-error {{scalar operand type has greater rank than the type of the vector element. ('float2' (vector of 2 'float' values) and 'double')}}
 
 kernel void foo2(float2 in, global float2 *out) { *out = 0.5 + in;} // expected-error {{scalar operand type has greater rank than the type of the vector element. ('double' and 'float2' (vector of 2 'float' values))}}
 
 kernel void foo3(float2 in, global float2 *out) { *out = 0.5f + in;}
 
 kernel void foo4(long2 in, global long2 *out) { *out = 5 + in;}
 
 kernel void foo5(float2 in, global float2 *out) {
     float* f;
-    *out = f + in; // expected-error{{cannot convert between vector and non-scalar values ('float *' and 'float2' (vector of 2 'float' values))}}
+    *out = f + in; // expected-error{{cannot convert between vector and non-scalar values ('__private float *' and 'float2' (vector of 2 'float' values))}}
 }
 
 kernel void foo6(int2 in, global int2 *out) {
     int* f;
-    *out = f + in; // expected-error{{cannot convert between vector and non-scalar values ('int *' and 'int2' (vector of 2 'int' values))}}
+    *out = f + in; // expected-error{{cannot convert between vector and non-scalar values ('__private int *' and 'int2' (vector of 2 'int' values))}}
 }

clang/test/SemaOpenCL/as_type.cl

 // RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -finclude-default-header -o - -verify -fsyntax-only
 
 char3 f1(char16 x) {
-  return  __builtin_astype(x, char3); // expected-error{{invalid reinterpretation: sizes of 'char3' (vector of 3 'char' values) and 'char16' (vector of 16 'char' values) must match}}
+  return  __builtin_astype(x, char3); // expected-error{{invalid reinterpretation: sizes of 'char3' (vector of 3 'char' values) and '__private char16' (vector of 16 'char' values) must match}}
 }
 
 char16 f3(int x) {
-  return __builtin_astype(x, char16); // expected-error{{invalid reinterpretation: sizes of 'char16' (vector of 16 'char' values) and 'int' must match}}
+  return __builtin_astype(x, char16); // expected-error{{invalid reinterpretation: sizes of 'char16' (vector of 16 'char' values) and '__private int' must match}}
 }
 
 void foo() {
     char src = 1;
-    int dst = as_int(src); // expected-error{{invalid reinterpretation: sizes of 'int' and 'char' must match}}
+    int dst = as_int(src); // expected-error{{invalid reinterpretation: sizes of 'int' and '__private char' must match}}
 }

clang/test/SemaOpenCL/atomic-ops.cl

void f(atomic_int *i, const atomic_int *ci,
 
   __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)}}
 
   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 '__generic int *__private' 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 *__private' 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 '__generic int *__private' 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 *__private' 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);
 
   __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'}}
 }
 
 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;
-  (void)__opencl_atomic_load(Ap, s, memory_scope_work_group); // expected-error {{passing 'struct S' to parameter of incompatible type 'int'}}
+  (void)__opencl_atomic_load(Ap, s, memory_scope_work_group); // expected-error {{passing '__private struct S' to parameter of incompatible type 'int'}}
 
   (void)__opencl_atomic_load(Ap, memory_order_relaxed, memory_scope_work_group);
   (void)__opencl_atomic_load(Ap, memory_order_acquire, memory_scope_work_group);
   (void)__opencl_atomic_load(Ap, memory_order_consume, memory_scope_work_group); // expected-error {{use of undeclared identifier 'memory_order_consume'}}
   (void)__opencl_atomic_load(Ap, memory_order_release, memory_scope_work_group); // expected-warning {{memory order argument to atomic operation is invalid}}
   (void)__opencl_atomic_load(Ap, memory_order_acq_rel, memory_scope_work_group); // expected-warning {{memory order argument to atomic operation is invalid}}
   (void)__opencl_atomic_load(Ap, memory_order_seq_cst, memory_scope_work_group);
 
void synchscope_checks(atomic_int *Ap, int scope) {
   (void)__opencl_atomic_load(Ap, memory_order_relaxed, scope);
   (void)__opencl_atomic_load(Ap, memory_order_relaxed, 10);    //expected-error{{synchronization scope argument to atomic operation is invalid}}
 
   // non-integer memory scope is casted to integer type.
   float fscope;
   (void)__opencl_atomic_load(Ap, memory_order_relaxed, 1.0f);
   (void)__opencl_atomic_load(Ap, memory_order_relaxed, fscope);
   struct S s;
-  (void)__opencl_atomic_load(Ap, memory_order_relaxed, s); //expected-error{{passing 'struct S' to parameter of incompatible type 'int'}}
+  (void)__opencl_atomic_load(Ap, memory_order_relaxed, s); //expected-error{{passing '__private struct S' to parameter of incompatible type 'int'}}
 }
 
 void nullPointerWarning(atomic_int *Ap, int *p, int val) {
   // The 'expected' pointer shouldn't be NULL.
   (void)__opencl_atomic_compare_exchange_strong(Ap, (void *)0, val, memory_order_relaxed, memory_order_relaxed, memory_scope_work_group); // expected-warning {{null passed to a callee that requires a non-null argument}}
 }

clang/test/SemaOpenCL/cl20-device-side-enqueue.cl

#endif
 
   char c;
   enqueue_kernel(default_queue, flags, ndrange,
                  ^(local void *a, local void *b) {
                    return 0;
                  },
                  c, 1024L);
 #ifdef WCONV
-// expected-warning-re@-2{{implicit conversion changes signedness: 'char' to 'unsigned {{int|long}}'}}
+// expected-warning-re@-2{{implicit conversion changes signedness: '__private char' to 'unsigned {{int|long}}'}}
 #endif
 #define UINT_MAX 4294967295
 
   enqueue_kernel(default_queue, flags, ndrange,
                  ^(local void *a, local void *b) {
                    return 0;
                  },
                  sizeof(int), sizeof(int) * UINT_MAX);

clang/test/SemaOpenCL/clk_event_t.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0
 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++
 
 // Taken from opencl-c.h
 #define CLK_NULL_EVENT (__builtin_astype(((__SIZE_MAX__)), clk_event_t))
 
 global clk_event_t ce; // expected-error {{the '__global clk_event_t' type cannot be used to declare a program scope variable}}
 
 int clk_event_tests() {
   event_t e;
   clk_event_t ce1;
   clk_event_t ce2;
   clk_event_t ce3 = CLK_NULL_EVENT;
 
-  if (e == ce1) { // expected-error {{invalid operands to binary expression ('event_t' and 'clk_event_t')}}
+  if (e == ce1) { // expected-error {{invalid operands to binary expression ('__private event_t' and '__private clk_event_t')}}
     return 9;
   }
 
   if (ce1 != ce2) {
     return 1;
   }
   else if (ce1 == CLK_NULL_EVENT || ce2 != CLK_NULL_EVENT) {
     return 0;
   }
 
   return 2;
 }

clang/test/SemaOpenCL/event_t.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
 
-event_t glb_evt; // expected-error {{the 'event_t' type cannot be used to declare a program scope variable}} expected-error{{program scope variable must reside in constant address space}}
+event_t glb_evt; // expected-error {{the '__private event_t' type cannot be used to declare a program scope variable}} expected-error{{program scope variable must reside in constant address space}}
 
 constant struct evt_s {
   event_t evt; // expected-error {{the 'event_t' type cannot be used to declare a structure or union field}}
 } evt_str = {0};
 
 void foo(event_t evt); // expected-note {{passing argument to parameter 'evt' here}}
 
-void kernel ker(event_t argevt) { // expected-error {{'event_t' cannot be used as the type of a kernel parameter}}
+void kernel ker(event_t argevt) { // expected-error {{'__private event_t' cannot be used as the type of a kernel parameter}}
   event_t e;
   constant event_t const_evt; // expected-error {{the event_t type can only be used with __private address space qualifier}} expected-error{{variable in constant address space must be initialized}}
   foo(e);
   foo(0);
   foo(5); // expected-error {{passing 'int' to parameter of incompatible type 'event_t'}}
   foo((event_t)1); // expected-error {{cannot cast non-zero value '1' to 'event_t'}}
 }
 

clang/test/SemaOpenCL/extension-begin.cl

void test_f1(void) {
   g(0);
 }
 
 #pragma OPENCL EXTENSION my_ext : disable 
 void test_f2(void) {
   struct A test_A2; // expected-error {{use of type 'struct A' requires my_ext extension to be enabled}}
   const struct A test_A_local; // expected-error {{use of type 'struct A' requires my_ext extension to be enabled}}
   TypedefOfA test_typedef_A; // expected-error {{use of type 'TypedefOfA' (aka 'struct A') requires my_ext extension to be enabled}}
-  PointerOfA test_A_pointer; // expected-error {{use of type 'PointerOfA' (aka 'const struct A *') requires my_ext extension to be enabled}}
+  PointerOfA test_A_pointer; // expected-error {{use of type 'PointerOfA' (aka 'const __private struct A *') requires my_ext extension to be enabled}}
   f(); // expected-error {{use of declaration 'f' requires my_ext extension to be enabled}}
   g(0); // expected-error {{no matching function for call to 'g'}}
         // expected-note@extension-begin.h:18 {{candidate unavailable as it requires OpenCL extension 'my_ext' to be enabled}}
         // expected-note@extension-begin.h:23 {{candidate function not viable: requires 0 arguments, but 1 was provided}}
 }
 

clang/test/SemaOpenCL/half.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -Wno-unused-value -triple spir-unknown-unknown
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : disable
 constant float f = 1.0h; // expected-error{{half precision constant requires cl_khr_fp16}}
 
 half half_disabled(half *p, // expected-error{{declaring function return value of type 'half' is not allowed}}
-                   half h)  // expected-error{{declaring function parameter of type 'half' is not allowed}}
+                   half h)  // expected-error{{declaring function parameter of type '__private half' is not allowed}}
 {
-  half a[2]; // expected-error{{declaring variable of type 'half [2]' is not allowed}}
-  half b;    // expected-error{{declaring variable of type 'half' is not allowed}}
-  *p; // expected-error{{loading directly from pointer to type 'half' requires cl_khr_fp16. Use vector data load builtin functions instead}}
+  half a[2]; // expected-error{{declaring variable of type '__private half [2]' is not allowed}}
+  half b;    // expected-error{{declaring variable of type '__private half' is not allowed}}
+  *p; // expected-error{{loading directly from pointer to type '__private half' requires cl_khr_fp16. Use vector data load builtin functions instead}}
   *p = 0; // expected-error{{assigning directly to pointer to type 'half' requires cl_khr_fp16. Use vector data store builtin functions instead}}
-  p[1]; // expected-error{{loading directly from pointer to type 'half' requires cl_khr_fp16. Use vector data load builtin functions instead}}
+  p[1]; // expected-error{{loading directly from pointer to type '__private half' requires cl_khr_fp16. Use vector data load builtin functions instead}}
   p[1] = 0; // expected-error{{assigning directly to pointer to type 'half' requires cl_khr_fp16. Use vector data store builtin functions instead}}
 
   float c = 1.0f;
   b = (half) c;  // expected-error{{casting to type 'half' is not allowed}}
   c = (float) 1.0h;  // expected-error{{half precision constant requires cl_khr_fp16}}
   b = 1.0h; // expected-error{{half precision constant requires cl_khr_fp16}}
 
   half *allowed = &p[1];
   half *allowed2 = &*p;
   half *allowed3 = p + 1;
 
   return h;
 }
 
 kernel void half_disabled_kernel(global half *p,
-                                 half h);  // expected-error{{declaring function parameter of type 'half' is not allowed; did you forget * ?}}
+                                 half h);  // expected-error{{declaring function parameter of type '__private half' is not allowed; did you forget * ?}}
 
 // Exactly the same as above but with the cl_khr_fp16 extension enabled.
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 constant half a = 1.0h;
 half half_enabled(half *p, half h)
 {
   half a[2];
   half b;

clang/test/SemaOpenCL/images.cl

 // RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only
 
 void img2d_ro(read_only image2d_t); // expected-note 3{{passing argument to parameter here}}
 void img2d_wo(write_only image2d_t); // expected-note 2{{passing argument to parameter here}}
 void img2d_rw(read_write image2d_t); // expected-note 2{{passing argument to parameter here}}
 void img2d_default(image2d_t); // expected-note 2{{passing argument to parameter here}}
 
 void imgage_access_test(image2d_t img2dro, image3d_t img3dro) {
   img2d_ro(img2dro); // read_only = read_only
-  img2d_ro(img3dro); // expected-error{{passing '__read_only image3d_t' to parameter of incompatible type '__read_only image2d_t'}}
+  img2d_ro(img3dro); // expected-error{{passing '__private __read_only image3d_t' to parameter of incompatible type '__read_only image2d_t'}}
 }
 
 kernel void read_only_access_test(read_only image2d_t img) {
   img2d_ro(img); // read_only = read_only
-  img2d_wo(img); // expected-error {{passing '__read_only image2d_t' to parameter of incompatible type '__write_only image2d_t'}}
-  img2d_rw(img); // expected-error {{passing '__read_only image2d_t' to parameter of incompatible type '__read_write image2d_t'}}
+  img2d_wo(img); // expected-error {{passing '__private __read_only image2d_t' to parameter of incompatible type '__write_only image2d_t'}}
+  img2d_rw(img); // expected-error {{passing '__private __read_only image2d_t' to parameter of incompatible type '__read_write image2d_t'}}
   img2d_default(img); // read_only = read_only
 }
 
 kernel void write_only_access_test(write_only image2d_t img) {
-  img2d_ro(img); // expected-error {{passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
+  img2d_ro(img); // expected-error {{passing '__private __write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
   img2d_wo(img); // write_only = write_only
-  img2d_rw(img); // expected-error {{passing '__write_only image2d_t' to parameter of incompatible type '__read_write image2d_t'}}
-  img2d_default(img); // expected-error {{passing '__write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
+  img2d_rw(img); // expected-error {{passing '__private __write_only image2d_t' to parameter of incompatible type '__read_write image2d_t'}}
+  img2d_default(img); // expected-error {{passing '__private __write_only image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
 }
 
 kernel void read_write_access_test(read_write image2d_t img) {
-  img2d_ro(img);  // expected-error {{passing '__read_write image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
-  img2d_wo(img); // expected-error {{passing '__read_write image2d_t' to parameter of incompatible type '__write_only image2d_t'}}
+  img2d_ro(img);  // expected-error {{passing '__private __read_write image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
+  img2d_wo(img); // expected-error {{passing '__private __read_write image2d_t' to parameter of incompatible type '__write_only image2d_t'}}
   img2d_rw(img); //read_write = read_write
-  img2d_default(img); // expected-error {{passing '__read_write image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
+  img2d_default(img); // expected-error {{passing '__private __read_write image2d_t' to parameter of incompatible type '__read_only image2d_t'}}
 }

clang/test/SemaOpenCL/intel-subgroup-avc-ext-types.cl

 //   * In previous versions of the extension these macroses was defined as {0},
 //     so initialization with initializer list containing one integer equal to
 //     zero should also work
 
 struct st{};
 // negative test cases for initializers
 void foo(char c, float f, void* v, struct st ss) {
   intel_sub_group_avc_mce_payload_t payload_mce = 0; // No zero initializer for mce types
-  // expected-error@-1 {{initializing 'intel_sub_group_avc_mce_payload_t' with an expression of incompatible type 'int'}}
+  // expected-error@-1 {{initializing '__private intel_sub_group_avc_mce_payload_t' with an expression of incompatible type 'int'}}
   intel_sub_group_avc_ime_payload_t payload_ime = 1; // No literal initializer for *payload_t types
-  // expected-error@-1 {{initializing 'intel_sub_group_avc_ime_payload_t' with an expression of incompatible type 'int'}}
+  // expected-error@-1 {{initializing '__private intel_sub_group_avc_ime_payload_t' with an expression of incompatible type 'int'}}
   intel_sub_group_avc_ref_payload_t payload_ref = f;
-  // expected-error@-1 {{initializing 'intel_sub_group_avc_ref_payload_t' with an expression of incompatible type 'float'}}
+  // expected-error@-1 {{initializing '__private intel_sub_group_avc_ref_payload_t' with an expression of incompatible type '__private float'}}
   intel_sub_group_avc_sic_payload_t payload_sic = ss;
-  // expected-error@-1 {{initializing 'intel_sub_group_avc_sic_payload_t' with an expression of incompatible type 'struct st'}}
+  // expected-error@-1 {{initializing '__private intel_sub_group_avc_sic_payload_t' with an expression of incompatible type '__private struct st'}}
 
   intel_sub_group_avc_mce_result_t result_mce = 0; // No zero initializer for mce types
-  // expected-error@-1 {{initializing 'intel_sub_group_avc_mce_result_t' with an expression of incompatible type 'int'}}
+  // expected-error@-1 {{initializing '__private intel_sub_group_avc_mce_result_t' with an expression of incompatible type 'int'}}
   intel_sub_group_avc_ime_result_t result_ime = 1; // No literal initializer for *result_t types
-  // expected-error@-1 {{initializing 'intel_sub_group_avc_ime_result_t' with an expression of incompatible type 'int'}}
+  // expected-error@-1 {{initializing '__private intel_sub_group_avc_ime_result_t' with an expression of incompatible type 'int'}}
   intel_sub_group_avc_ref_result_t result_ref = f;
-  // expected-error@-1 {{initializing 'intel_sub_group_avc_ref_result_t' with an expression of incompatible type 'float'}}
+  // expected-error@-1 {{initializing '__private intel_sub_group_avc_ref_result_t' with an expression of incompatible type '__private float'}}
   intel_sub_group_avc_sic_result_t result_sic = ss;
-  // expected-error@-1 {{initializing 'intel_sub_group_avc_sic_result_t' with an expression of incompatible type 'struct st'}}
+  // expected-error@-1 {{initializing '__private intel_sub_group_avc_sic_result_t' with an expression of incompatible type '__private struct st'}}
 
   intel_sub_group_avc_ime_result_single_reference_streamout_t sstreamout = v;
-  // expected-error@-1 {{initializing 'intel_sub_group_avc_ime_result_single_reference_streamout_t' with an expression of incompatible type 'void *'}}
+  // expected-error@-1 {{initializing '__private intel_sub_group_avc_ime_result_single_reference_streamout_t' with an expression of incompatible type '__private void *__private'}}
 
   intel_sub_group_avc_ime_result_dual_reference_streamout_t dstreamin_list = {0x0, 0x1};
   // expected-warning@-1 {{excess elements in struct initializer}}
   intel_sub_group_avc_ime_dual_reference_streamin_t dstreamin_list2 = {};
   // expected-error@-1 {{scalar initializer cannot be empty}}
   intel_sub_group_avc_ime_single_reference_streamin_t dstreamin_list3 = {c};
-  // expected-error@-1 {{initializing 'intel_sub_group_avc_ime_single_reference_streamin_t' with an expression of incompatible type 'char'}}
+  // expected-error@-1 {{initializing '__private intel_sub_group_avc_ime_single_reference_streamin_t' with an expression of incompatible type '__private char'}}
   intel_sub_group_avc_ime_dual_reference_streamin_t dstreamin_list4 = {1};
-  // expected-error@-1 {{initializing 'intel_sub_group_avc_ime_dual_reference_streamin_t' with an expression of incompatible type 'int'}}
+  // expected-error@-1 {{initializing '__private intel_sub_group_avc_ime_dual_reference_streamin_t' with an expression of incompatible type 'int'}}
 }
 
 // negative tests for initializers and assignment
 void far() {
   intel_sub_group_avc_mce_payload_t payload_mce;
   intel_sub_group_avc_mce_payload_t payload_mce2 = payload_mce;
 
   intel_sub_group_avc_ime_payload_t payload_ime;
   intel_sub_group_avc_ref_payload_t payload_ref = payload_ime;
-  // expected-error@-1 {{initializing 'intel_sub_group_avc_ref_payload_t' with an expression of incompatible type 'intel_sub_group_avc_ime_payload_t'}}
+  // expected-error@-1 {{initializing '__private intel_sub_group_avc_ref_payload_t' with an expression of incompatible type '__private intel_sub_group_avc_ime_payload_t'}}
 
   intel_sub_group_avc_sic_result_t result_sic;
   intel_sub_group_avc_ime_result_t result_ime;
   result_sic = result_ime;
-  // expected-error@-1 {{assigning to 'intel_sub_group_avc_sic_result_t' from incompatible type 'intel_sub_group_avc_ime_result_t'}}
+  // expected-error@-1 {{assigning to '__private intel_sub_group_avc_sic_result_t' from incompatible type '__private intel_sub_group_avc_ime_result_t'}}
 }
 
 // Using 0x0 directly allows us not to include opencl-c.h header and not to
 // redefine all of these CLK_AVC_*_INTITIALIZE_INTEL macro. '0x0' value must
 // be in sync with ones defined in opencl-c.h
 
 // positive test cases
 void bar() {

clang/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 (__generic ^const __private)(void)' and 'int (__generic ^const __private)(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() {

 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 (__generic ^const)(__private 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 'bl2_t' (aka 'int (__generic ^const)(int)') is invalid in OpenCL}}
+void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type 'bl2_t' (aka 'int (__generic ^const)(__private int)') is invalid in OpenCL}}
   bl2_t bl = ^(int i) {
     return 1;
   };
-  bl2_t *p; // expected-error {{pointer to type 'bl2_t' (aka 'int (__generic ^const)(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 (__generic ^const)(__private int)') is invalid in OpenCL}}
+  *bl;      // expected-error {{invalid argument type '__private bl2_t' (aka 'int (__generic ^const __private)(__private int)') to unary expression}}
+  &bl;      // expected-error {{invalid argument type '__private bl2_t' (aka 'int (__generic ^const __private)(__private 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}}

clang/test/SemaOpenCL/invalid-image.cl

 // RUN: %clang_cc1 -verify -cl-std=clc++ %s
 // RUN: %clang_cc1 -verify %s
 // RUN: %clang_cc1 -verify -D=ATTR_TEST -fms-compatibility %s
 
 void test1(image1d_t *i) {} // expected-error-re{{pointer to type '{{__generic __read_only|__read_only}} image1d_t' is invalid in OpenCL}}
 
 void test2(image1d_t i) {
-  image1d_t ti;            // expected-error{{type '__read_only image1d_t' can only be used as a function parameter}}
+  image1d_t ti;            // expected-error{{type '__private __read_only image1d_t' can only be used as a function parameter}}
   image1d_t ai[] = {i, i}; // expected-error{{array of '__read_only image1d_t' type is invalid in OpenCL}}
-  i=i; // expected-error{{invalid operands to binary expression ('__read_only image1d_t' and '__read_only image1d_t')}}
-  i+1; // expected-error{{invalid operands to binary expression ('__read_only image1d_t' and 'int')}}
-  &i; // expected-error{{invalid argument type '__read_only image1d_t' to unary expression}}
-  *i; // expected-error{{invalid argument type '__read_only image1d_t' to unary expression}}
+  i=i; // expected-error{{invalid operands to binary expression ('__private __read_only image1d_t' and '__private __read_only image1d_t')}}
+  i+1; // expected-error{{invalid operands to binary expression ('__private __read_only image1d_t' and 'int')}}
+  &i; // expected-error{{invalid argument type '__private __read_only image1d_t' to unary expression}}
+  *i; // expected-error{{invalid argument type '__private __read_only image1d_t' to unary expression}}
 }
 
 image1d_t test3() {} // expected-error{{declaring function return value of type '__read_only image1d_t' is not allowed}}
 
 #ifdef ATTR_TEST
 // Test case for an infinite loop bug.
 kernel void foob(read_only __ptr32  image2d_t i) { } // expected-error{{'__ptr32' attribute only applies to pointer arguments}}
 #endif

clang/test/SemaOpenCL/invalid-kernel-parameters.cl

 // RUN: %clang_cc1 -fsyntax-only -verify %s -triple spir-unknown-unknown
 
-kernel void half_arg(half x) { } // expected-error{{declaring function parameter of type 'half' is not allowed; did you forget * ?}}
+kernel void half_arg(half x) { } // expected-error{{declaring function parameter of type '__private half' is not allowed; did you forget * ?}}
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
 
 // Disallowed: parameters with type
 // bool, half, size_t, ptrdiff_t, intptr_t, and uintptr_t
 // or a struct / union with any of these types in them
 
-typedef __SIZE_TYPE__ size_t; // expected-note{{'size_t' (aka 'unsigned int') declared here}}
+typedef __SIZE_TYPE__ size_t; // expected-note{{'__private size_t' (aka '__private unsigned int') declared here}}
                               // expected-note@-1{{'size_t' (aka 'unsigned int') declared here}}
-typedef __PTRDIFF_TYPE__ ptrdiff_t; // expected-note{{'ptrdiff_t' (aka 'int') declared here}}
-typedef __INTPTR_TYPE__ intptr_t; // expected-note{{'intptr_t' (aka 'int') declared here}}
-typedef __UINTPTR_TYPE__ uintptr_t; // expected-note{{'uintptr_t' (aka 'unsigned int') declared here}}
+typedef __PTRDIFF_TYPE__ ptrdiff_t; // expected-note{{'__private ptrdiff_t' (aka '__private int') declared here}}
+typedef __INTPTR_TYPE__ intptr_t; // expected-note{{'__private intptr_t' (aka '__private int') declared here}}
+typedef __UINTPTR_TYPE__ uintptr_t; // expected-note{{'__private uintptr_t' (aka '__private unsigned int') declared here}}
 
-kernel void size_t_arg(size_t x) {} // expected-error{{'size_t' (aka 'unsigned int') cannot be used as the type of a kernel parameter}}
+kernel void size_t_arg(size_t x) {} // expected-error{{'__private size_t' (aka '__private unsigned int') cannot be used as the type of a kernel parameter}}
 
-kernel void ptrdiff_t_arg(ptrdiff_t x) {} // expected-error{{'ptrdiff_t' (aka 'int') cannot be used as the type of a kernel parameter}}
+kernel void ptrdiff_t_arg(ptrdiff_t x) {} // expected-error{{'__private ptrdiff_t' (aka '__private int') cannot be used as the type of a kernel parameter}}
 
-kernel void intptr_t_arg(intptr_t x) {} // expected-error{{'intptr_t' (aka 'int') cannot be used as the type of a kernel parameter}}
+kernel void intptr_t_arg(intptr_t x) {} // expected-error{{'__private intptr_t' (aka '__private int') cannot be used as the type of a kernel parameter}}
 
-kernel void uintptr_t_arg(uintptr_t x) {} // expected-error{{'uintptr_t' (aka 'unsigned int') cannot be used as the type of a kernel parameter}}
+kernel void uintptr_t_arg(uintptr_t x) {} // expected-error{{'__private uintptr_t' (aka '__private unsigned int') cannot be used as the type of a kernel parameter}}
 
 typedef size_t size_ty;
 struct SizeTStruct { // expected-note{{within field of type 'SizeTStruct' declared here}}
   size_ty s; // expected-note{{field of illegal type 'size_ty' (aka 'unsigned int') declared here}}
 };
-kernel void size_t_struct_arg(struct SizeTStruct x) {} // expected-error{{'struct SizeTStruct' cannot be used as the type of a kernel parameter}}
+kernel void size_t_struct_arg(struct SizeTStruct x) {} // expected-error{{'__private struct SizeTStruct' cannot be used as the type of a kernel parameter}}
 
 union SizeTUnion { // expected-note{{within field of type 'SizeTUnion' declared here}}
   size_t s; // expected-note{{field of illegal type 'size_t' (aka 'unsigned int') declared here}}
   float f;
 };
-kernel void size_t_union_arg(union SizeTUnion x) {} // expected-error{{'union SizeTUnion' cannot be used as the type of a kernel parameter}}
+kernel void size_t_union_arg(union SizeTUnion x) {} // expected-error{{'__private union SizeTUnion' cannot be used as the type of a kernel parameter}}
 
 typedef size_t s_ty; // expected-note{{'s_ty' (aka 'unsigned int') declared here}}
-typedef s_ty ss_ty; // expected-note{{'ss_ty' (aka 'unsigned int') declared here}}
-kernel void typedef_to_size_t(ss_ty s) {} // expected-error{{'ss_ty' (aka 'unsigned int') cannot be used as the type of a kernel parameter}}
+typedef s_ty ss_ty; // expected-note{{'__private ss_ty' (aka '__private unsigned int') declared here}}
+kernel void typedef_to_size_t(ss_ty s) {} // expected-error{{'__private ss_ty' (aka '__private unsigned int') cannot be used as the type of a kernel parameter}}
 
-kernel void bool_arg(bool x) { } // expected-error{{'bool' cannot be used as the type of a kernel parameter}}
+kernel void bool_arg(bool x) { } // expected-error{{'__private bool' cannot be used as the type of a kernel parameter}}
 
 // half kernel argument is allowed when cl_khr_fp16 is enabled.
 kernel void half_arg(half x) { }
 
 typedef struct ContainsBool // expected-note{{within field of type 'ContainsBool' declared here}}
 {
   bool x; // expected-note{{field of illegal type 'bool' declared here}}
 } ContainsBool;
 
-kernel void bool_in_struct_arg(ContainsBool x) { } // expected-error{{'ContainsBool' (aka 'struct ContainsBool') cannot be used as the type of a kernel parameter}}
+kernel void bool_in_struct_arg(ContainsBool x) { } // expected-error{{'__private ContainsBool' (aka '__private struct ContainsBool') cannot be used as the type of a kernel parameter}}
 
 
 
 typedef struct FooImage2D // expected-note{{within field of type 'FooImage2D' declared here}}
 {
   // TODO: Clean up needed - we don't really need to check for image, event, etc
   // as a note here any longer.
   // They are diagnosed as an error for all struct fields (OpenCL v1.2 s6.9b,r).
   image2d_t imageField; // expected-note{{field of illegal type '__read_only image2d_t' declared here}} expected-error{{the '__read_only image2d_t' type cannot be used to declare a structure or union field}}
 } FooImage2D;
 
 kernel void image_in_struct_arg(FooImage2D arg) { } // expected-error{{struct kernel parameters may not contain pointers}}
 
 typedef struct Foo // expected-note{{within field of type 'Foo' declared here}}
 {
-  int* ptrField; // expected-note{{field of illegal pointer type 'int *' declared here}}
+  int* ptrField; // expected-note{{field of illegal pointer type '__private int *' declared here}}
 } Foo;
 
 kernel void pointer_in_struct_arg(Foo arg) { } // expected-error{{struct kernel parameters may not contain pointers}}
 
 typedef union FooUnion // expected-note{{within field of type 'FooUnion' declared here}}
 {
-  int* ptrField; // expected-note{{field of illegal pointer type 'int *' declared here}}
+  int* ptrField; // expected-note{{field of illegal pointer type '__private int *' declared here}}
 } FooUnion;
 
 kernel void pointer_in_union_arg(FooUnion arg) { }// expected-error{{union kernel parameters may not contain pointers}}
 
 typedef struct NestedPointer // expected-note 2 {{within field of type 'NestedPointer' declared here}}
 {
   int x;
   struct InnerNestedPointer
   {
-    int* ptrField; // expected-note 3 {{field of illegal pointer type 'int *' declared here}}
+    int* ptrField; // expected-note 3 {{field of illegal pointer type '__private int *' declared here}}
   } inner; // expected-note 3 {{within field of type 'struct InnerNestedPointer' declared here}}
 } NestedPointer;
 
 kernel void pointer_in_nested_struct_arg(NestedPointer arg) { }// expected-error{{struct kernel parameters may not contain pointers}}
 
 struct NestedPointerComplex // expected-note{{within field of type 'NestedPointerComplex' declared here}}
 {
   int foo;
   float bar;
 
   struct InnerNestedPointerComplex
   {
     int innerFoo;
-    int* innerPtrField; // expected-note{{field of illegal pointer type 'int *' declared here}}
+    int* innerPtrField; // expected-note{{field of illegal pointer type '__private int *' declared here}}
   } inner; // expected-note{{within field of type 'struct InnerNestedPointerComplex' declared here}}
 
   float y;
   float z[4];
 };
 
 kernel void pointer_in_nested_struct_arg_complex(struct NestedPointerComplex arg) { }// expected-error{{struct kernel parameters may not contain pointers}}
 
 typedef struct NestedBool // expected-note 2 {{within field of type 'NestedBool' declared here}}
 {
   int x;
   struct InnerNestedBool
   {
     bool boolField; // expected-note 2 {{field of illegal type 'bool' declared here}}
   } inner; // expected-note 2 {{within field of type 'struct InnerNestedBool' declared here}}
 } NestedBool;
 
-kernel void bool_in_nested_struct_arg(NestedBool arg) { } // expected-error{{'NestedBool' (aka 'struct NestedBool') cannot be used as the type of a kernel parameter}}
+kernel void bool_in_nested_struct_arg(NestedBool arg) { } // expected-error{{'__private NestedBool' (aka '__private struct NestedBool') cannot be used as the type of a kernel parameter}}
 
 // Warning emitted again for argument used in other kernel
-kernel void bool_in_nested_struct_arg_again(NestedBool arg) { } // expected-error{{'NestedBool' (aka 'struct NestedBool') cannot be used as the type of a kernel parameter}}
+kernel void bool_in_nested_struct_arg_again(NestedBool arg) { } // expected-error{{'__private NestedBool' (aka '__private struct NestedBool') cannot be used as the type of a kernel parameter}}
 
 
 // Check for note with a struct not defined inside the struct
 typedef struct NestedBool2Inner
 {
   bool boolField; // expected-note{{field of illegal type 'bool' declared here}}
 } NestedBool2Inner;
 
 typedef struct NestedBool2 // expected-note{{within field of type 'NestedBool2' declared here}}
 {
   int x;
   NestedBool2Inner inner; // expected-note{{within field of type 'NestedBool2Inner' (aka 'struct NestedBool2Inner') declared here}}
 } NestedBool2;
 
-kernel void bool_in_nested_struct_2_arg(NestedBool2 arg) { } // expected-error{{'NestedBool2' (aka 'struct NestedBool2') cannot be used as the type of a kernel parameter}}
+kernel void bool_in_nested_struct_2_arg(NestedBool2 arg) { } // expected-error{{'__private NestedBool2' (aka '__private struct NestedBool2') cannot be used as the type of a kernel parameter}}
 
 
 struct InnerInner
 {
   int* foo;
   bool x;
 };
 
struct AlsoUser // expected-note{{within field of type 'AlsoUser' declared here}}
   struct Valid valid2;
   struct NestedPointer aaaa; // expected-note{{within field of type 'struct NestedPointer' declared here}}
 };
 
 kernel void pointer_in_nested_struct_arg_2(struct Valid valid, struct NestedPointer arg, struct AlsoUser also) { } // expected-error 2 {{struct kernel parameters may not contain pointers}}
 
 struct ArrayOfPtr // expected-note{{within field of type 'ArrayOfPtr' declared here}}
 {
-  float *arr[3]; // expected-note{{field of illegal type 'float *[3]' declared here}}
-                 // expected-note@-1{{field of illegal type 'float *[3]' declared here}}
+  float *arr[3]; // expected-note{{field of illegal type '__private float *[3]' declared here}}
+                 // expected-note@-1{{field of illegal type '__private float *[3]' declared here}}
 };
 kernel void array_of_ptr(struct ArrayOfPtr arr) {} // expected-error{{struct kernel parameters may not contain pointers}}
 
 struct ArrayOfStruct // expected-note{{within field of type 'ArrayOfStruct' declared here}}
 {
   struct ArrayOfPtr arr[3]; // expected-note{{within field of type 'struct ArrayOfPtr [3]' declared here}}
 };
 kernel void array_of_struct(struct ArrayOfStruct arr) {} // expected-error{{struct kernel parameters may not contain pointers}}

clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0 -cl-ext=+cl_khr_subgroups
 
 #pragma OPENCL EXTENSION cl_khr_subgroups : enable
 
 void test1(read_only pipe int p, global int* ptr){
   int tmp;
   reserve_id_t rid;
 
   // read/write_pipe
   read_pipe(p, &tmp);
   read_pipe(p, ptr);
   read_pipe(tmp, p);    // expected-error {{first argument to 'read_pipe' must be a pipe type}}
   read_pipe(p);   // expected-error {{invalid number of arguments to function: 'read_pipe'}}
   read_pipe(p, rid, tmp, ptr);
-  read_pipe(p, tmp, tmp, ptr);   // expected-error {{invalid argument type to function 'read_pipe' (expecting 'reserve_id_t' having 'int')}}
-  read_pipe(p, rid, rid, ptr);   // expected-error {{invalid argument type to function 'read_pipe' (expecting 'unsigned int' having 'reserve_id_t')}}
-  read_pipe(p, tmp);   // expected-error {{invalid argument type to function 'read_pipe' (expecting 'int *' having 'int')}}
+  read_pipe(p, tmp, tmp, ptr);   // expected-error {{invalid argument type to function 'read_pipe' (expecting 'reserve_id_t' having '__private int')}}
+  read_pipe(p, rid, rid, ptr);   // expected-error {{invalid argument type to function 'read_pipe' (expecting 'unsigned int' having '__private reserve_id_t')}}
+  read_pipe(p, tmp);   // expected-error {{invalid argument type to function 'read_pipe' (expecting 'int *' having '__private int')}}
   write_pipe(p, ptr);    // expected-error {{invalid pipe access modifier (expecting write_only)}}
   write_pipe(p, rid, tmp, ptr);    // expected-error {{invalid pipe access modifier (expecting write_only)}}
 
   // reserve_read/write_pipe
   reserve_read_pipe(p, tmp);
-  reserve_read_pipe(p, ptr);    // expected-error{{invalid argument type to function 'reserve_read_pipe' (expecting 'unsigned int' having '__global int *')}}
+  reserve_read_pipe(p, ptr);    // expected-error{{invalid argument type to function 'reserve_read_pipe' (expecting 'unsigned int' having '__global int *__private')}}
   work_group_reserve_read_pipe(tmp, tmp);    // expected-error{{first argument to 'work_group_reserve_read_pipe' must be a pipe type}}
   sub_group_reserve_write_pipe(p, tmp);    // expected-error{{invalid pipe access modifier (expecting write_only)}}
 
   // commit_read/write_pipe
   commit_read_pipe(p, rid);
   commit_read_pipe(tmp, rid);    // expected-error{{first argument to 'commit_read_pipe' must be a pipe type}}
-  work_group_commit_read_pipe(p, tmp);    // expected-error{{invalid argument type to function 'work_group_commit_read_pipe' (expecting 'reserve_id_t' having 'int')}}
+  work_group_commit_read_pipe(p, tmp);    // expected-error{{invalid argument type to function 'work_group_commit_read_pipe' (expecting 'reserve_id_t' having '__private int')}}
   sub_group_commit_write_pipe(p, tmp);    // expected-error{{invalid pipe access modifier (expecting write_only)}}
 }
 
 void test2(write_only pipe int p, global int* ptr){
   int tmp;
   reserve_id_t rid;
 
   // read/write_pipe
   write_pipe(p, &tmp);
   write_pipe(p, ptr);
   write_pipe(tmp, p);    // expected-error {{first argument to 'write_pipe' must be a pipe type}}
   write_pipe(p);   // expected-error {{invalid number of arguments to function: 'write_pipe'}}
   write_pipe(p, rid, tmp, ptr);
-  write_pipe(p, tmp, tmp, ptr);   // expected-error {{invalid argument type to function 'write_pipe' (expecting 'reserve_id_t' having 'int')}}
-  write_pipe(p, rid, rid, ptr);   // expected-error {{invalid argument type to function 'write_pipe' (expecting 'unsigned int' having 'reserve_id_t')}}
-  write_pipe(p, tmp);   // expected-error {{invalid argument type to function 'write_pipe' (expecting 'int *' having 'int')}}
+  write_pipe(p, tmp, tmp, ptr);   // expected-error {{invalid argument type to function 'write_pipe' (expecting 'reserve_id_t' having '__private int')}}
+  write_pipe(p, rid, rid, ptr);   // expected-error {{invalid argument type to function 'write_pipe' (expecting 'unsigned int' having '__private reserve_id_t')}}
+  write_pipe(p, tmp);   // expected-error {{invalid argument type to function 'write_pipe' (expecting 'int *' having '__private int')}}
   read_pipe(p, ptr);    // expected-error {{invalid pipe access modifier (expecting read_only)}}
   read_pipe(p, rid, tmp, ptr);    // expected-error {{invalid pipe access modifier (expecting read_only)}}
 
   // reserve_read/write_pipe
   reserve_write_pipe(p, tmp);
-  reserve_write_pipe(p, ptr);    // expected-error{{invalid argument type to function 'reserve_write_pipe' (expecting 'unsigned int' having '__global int *')}}
+  reserve_write_pipe(p, ptr);    // expected-error{{invalid argument type to function 'reserve_write_pipe' (expecting 'unsigned int' having '__global int *__private')}}
   work_group_reserve_write_pipe(tmp, tmp);    // expected-error{{first argument to 'work_group_reserve_write_pipe' must be a pipe type}}
   sub_group_reserve_read_pipe(p, tmp);    // expected-error{{invalid pipe access modifier (expecting read_only)}}
 
   // commit_read/write_pipe
   commit_write_pipe(p, rid);
   commit_write_pipe(tmp, rid);    // expected-error{{first argument to 'commit_write_pipe' must be a pipe type}}
-  work_group_commit_write_pipe(p, tmp);    // expected-error{{invalid argument type to function 'work_group_commit_write_pipe' (expecting 'reserve_id_t' having 'int')}}
+  work_group_commit_write_pipe(p, tmp);    // expected-error{{invalid argument type to function 'work_group_commit_write_pipe' (expecting 'reserve_id_t' having '__private int')}}
   sub_group_commit_read_pipe(p, tmp);    // expected-error{{invalid pipe access modifier (expecting read_only)}}
 }
 
 void test3(){
   int tmp;
   get_pipe_num_packets(tmp);    // expected-error {{first argument to 'get_pipe_num_packets' must be a pipe type}}
   get_pipe_max_packets(tmp);    // expected-error {{first argument to 'get_pipe_max_packets' must be a pipe type}}
 }

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

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0
 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++
 
 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-re{{type '__global write_only pipe int ({{(void)?}})' can only be used as a function parameter in OpenCL}} expected-error{{'write_only' attribute only applies to parameters and typedefs}}
 
-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}}
+kernel void test_invalid_reserved_id(reserve_id_t ID) { // expected-error {{'__private 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}}
 }
 void test3(int pipe p) {// expected-error {{cannot combine with previous 'int' declaration specifier}}
 }
 void test4() {
-  pipe int p; // expected-error {{type 'read_only pipe int' can only be used as a function parameter}}
+  pipe int p; // expected-error {{type '__private read_only pipe int' can only be used as a function parameter}}
   //TODO: fix parsing of this pipe int (*p);
 }
 
 void test5(pipe int p) {
-  p+p; // expected-error{{invalid operands to binary expression ('read_only pipe int' and 'read_only pipe int')}}
-  p=p; // expected-error{{invalid operands to binary expression ('read_only pipe int' and 'read_only pipe int')}}
-  &p; // expected-error{{invalid argument type 'read_only pipe int' to unary expression}}
-  *p; // expected-error{{invalid argument type 'read_only pipe int' to unary expression}}
+  p+p; // expected-error{{invalid operands to binary expression ('__private read_only pipe int' and '__private read_only pipe int')}}
+  p=p; // expected-error{{invalid operands to binary expression ('__private read_only pipe int' and '__private read_only pipe int')}}
+  &p; // expected-error{{invalid argument type '__private read_only pipe int' to unary expression}}
+  *p; // expected-error{{invalid argument type '__private read_only pipe int' to unary expression}}
 }
 
 typedef pipe int pipe_int_t;
 pipe_int_t test6() {} // expected-error{{declaring function return value of type 'pipe_int_t' (aka 'read_only pipe int') is not allowed}}
 
 bool test_id_comprision(void) {
   reserve_id_t id1, id2;
-  return (id1 == id2);          // expected-error {{invalid operands to binary expression ('reserve_id_t' and 'reserve_id_t')}}
+  return (id1 == id2);          // expected-error {{invalid operands to binary expression ('__private reserve_id_t' and '__private reserve_id_t')}}
 }
 
 // Tests ASTContext::mergeTypes rejects this.
 #ifndef __OPENCL_CPP_VERSION__
 int f(pipe int x, int y); // expected-note {{previous declaration is here}}
 int f(x, y) // expected-error {{conflicting types for 'f}}
 pipe short x;
 int y;
 {
     return y;
 }
 #endif

clang/test/SemaOpenCL/null_literal.cl

 // RUN: %clang_cc1 -verify %s
 // RUN: %clang_cc1 -cl-std=CL2.0 -DCL20 -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}}
+constant int* ptr4 = (global void*)0; // expected-error{{initializing '__constant int *__private' 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;
 #endif
 
-global int* ptr6 = (local void*)0; // expected-error{{initializing '__global int *' with an expression of type '__local void *' changes address space of pointer}}
+global int* ptr6 = (local void*)0; // expected-error{{initializing '__global int *__private' 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;
 
 }

clang/test/SemaOpenCL/null_queue.cl

 // RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only
 extern queue_t get_default_queue();
 
 void queue_arg(queue_t); // expected-note {{passing argument to parameter here}}
 
 void init() {
-  queue_t q1 = 1; // expected-error{{initializing 'queue_t' with an expression of incompatible type 'int'}}
+  queue_t q1 = 1; // expected-error{{initializing '__private queue_t' with an expression of incompatible type 'int'}}
   queue_t q = 0;
 }
 
 void assign() {
   queue_t q2, q3;
-  q2 = 5; // expected-error{{assigning to 'queue_t' from incompatible type 'int'}}
+  q2 = 5; // expected-error{{assigning to '__private queue_t' from incompatible type 'int'}}
   q3 = 0;
   q2 = q3 = 0;
 }
 
 bool compare() {
   queue_t q4, q5;
   return 1 == get_default_queue() && // expected-error{{invalid operands to binary expression ('int' and 'queue_t')}}
          get_default_queue() == 1 && // expected-error{{invalid operands to binary expression ('queue_t' and 'int')}}
        q4 == q5 &&
        q4 != 0 &&
-       q4 != 0.0f; // expected-error{{invalid operands to binary expression ('queue_t' and 'float')}}
+       q4 != 0.0f; // expected-error{{invalid operands to binary expression ('__private queue_t' and 'float')}}
 }
 
 void call() {
   queue_arg(5); // expected-error {{passing 'int' to parameter of incompatible type 'queue_t'}}
   queue_arg(0);
 }

clang/test/SemaOpenCL/predefined-expr.cl

 // RUN: %clang_cc1 %s -verify
 // RUN: %clang_cc1 %s -verify -cl-std=CL2.0
 
 void f() {
-  char *f1 = __func__;          //expected-error-re{{initializing '{{__generic char|char}} *' with an expression of type 'const __constant char *' changes address space of pointer}}
-  constant char *f2 = __func__; //expected-warning{{initializing '__constant char *' with an expression of type 'const __constant char [2]' discards qualifiers}}
+  char *f1 = __func__;          //expected-error-re{{initializing '{{__generic|__private}} char *__private' with an expression of type 'const __constant char *' changes address space of pointer}}
+  constant char *f2 = __func__; //expected-warning{{initializing '__constant char *__private' with an expression of type 'const __constant char [2]' discards qualifiers}}
   constant const char *f3 = __func__;
 }

clang/test/SemaOpenCL/queue_t_overload.cl

 // RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only
 
-void __attribute__((overloadable)) foo(queue_t, __local char *); // expected-note {{candidate function not viable: no known conversion from 'int' to 'queue_t' for 1st argument}} // expected-note {{candidate function}}
-void __attribute__((overloadable)) foo(queue_t, __local float *); // expected-note {{candidate function not viable: no known conversion from 'int' to 'queue_t' for 1st argument}} // expected-note {{candidate function}}
+void __attribute__((overloadable)) foo(queue_t, __local char *); // expected-note {{candidate function not viable: no known conversion from 'int' to '__private queue_t' for 1st argument}} // expected-note {{candidate function}}
+void __attribute__((overloadable)) foo(queue_t, __local float *); // expected-note {{candidate function not viable: no known conversion from 'int' to '__private queue_t' for 1st argument}} // expected-note {{candidate function}}
 
 void kernel ker(__local char *src1, __local float *src2, __global int *src3) {
   queue_t q;
   foo(q, src1);
   foo(0, src2);
   foo(q, src3); // expected-error {{no matching function for call to 'foo'}}
   foo(1, src3); // expected-error {{no matching function for call to 'foo'}}
 }

clang/test/SemaOpenCL/shifts.cl

void ptest07() {
   char3 w = {1,2,3};
 
   v <<= w;
 }
 
 // ** Negative tests **
 
 char2 ntest01(char c, char2 s) {
-  return c << s; // expected-error {{requested shift is a vector of type 'char2' (vector of 2 'char' values) but the first operand is not a vector ('char')}}
+  return c << s; // expected-error {{requested shift is a vector of type '__private char2' (vector of 2 'char' values) but the first operand is not a vector ('__private char')}}
 }
 
 char3 ntest02(char3 c, char2 s) {
   return c << s; // expected-error {{vector operands do not have the same number of elements ('char3' (vector of 3 'char' values) and 'char2' (vector of 2 'char' values))}}
 }
 
 float2 ntest03(float2 c, char s) {
   return c << s; // expected-error {{used type 'float2' (vector of 2 'float' values) where integer is required}}
 }
 
 int2 ntest04(int2 c, float s) {
   return c << s; // expected-error {{used type 'float' where integer is required}}
 }

clang/test/SemaOpenCL/to_addr_builtin.cl

 // RUN: %clang_cc1 -verify -fsyntax-only %s
 // RUN: %clang_cc1 -Wconversion -verify -fsyntax-only -cl-std=CL2.0 %s
 
 void test(void) {
   global int *glob;
   local int *loc;
   constant int *con;
   private int *priv;
   global float *glob_wrong_ty;
   typedef constant int const_int_ty;
   const_int_ty *con_typedef;
 
   glob = to_global(glob, loc);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
   // expected-error@-2{{implicit declaration of function 'to_global' is invalid in OpenCL}}
-  // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
   // expected-error@-5{{invalid number of arguments to function: 'to_global'}}
 #endif
 
   int x;
   glob = to_global(x);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
   // expected-error@-4{{invalid argument x to function: 'to_global', expecting a generic pointer argument}}
 #endif
 
   glob = to_global(con);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
   // expected-error@-4{{invalid argument con to function: 'to_global', expecting a generic pointer argument}}
 #endif
 
   glob = to_global(con_typedef);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
   // expected-error@-4{{invalid argument con_typedef to function: 'to_global', expecting a generic pointer argument}}
 #endif
 
   loc = to_global(glob);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__local int *' from 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__local int *__private' from 'int'}}
 #else
-  // expected-error@-4{{assigning '__global int *' to '__local int *' changes address space of pointer}}
+  // expected-error@-4{{assigning '__global int *' to '__local int *__private' changes address space of pointer}}
   // expected-warning@-5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}
 #endif
 
   loc = to_private(glob);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
   // expected-error@-2{{implicit declaration of function 'to_private' is invalid in OpenCL}}
-  // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__local int *' from 'int'}}
+  // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__local int *__private' from 'int'}}
 #else
-  // expected-error@-5{{assigning 'int *' to '__local int *' changes address space of pointer}}
+  // expected-error@-5{{assigning '__private int *' to '__local int *__private' changes address space of pointer}}
   // expected-warning@-6{{passing non-generic address space pointer to to_private may cause dynamic conversion affecting performance}}
 #endif
 
   loc = to_local(glob);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
   // expected-error@-2{{implicit declaration of function 'to_local' is invalid in OpenCL}}
-  // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__local int *' from 'int'}}
+  // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__local int *__private' from 'int'}}
   // expected-note@-4{{did you mean 'to_global'}}
   // expected-note@13{{'to_global' declared here}}
 #else
   // expected-warning@-7{{passing non-generic address space pointer to to_local may cause dynamic conversion affecting performance}}
 #endif
 
   priv = to_global(glob);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion assigning to 'int *' from 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__private int *__private' from 'int'}}
 #else
-  // expected-error@-4{{assigning '__global int *' to 'int *' changes address space of pointer}}
+  // expected-error@-4{{assigning '__global int *' to '__private int *__private' changes address space of pointer}}
   // expected-warning@-5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}
 #endif
 
   priv = to_private(glob);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion assigning to 'int *' from 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__private int *__private' from 'int'}}
 #else
   // expected-warning@-4{{passing non-generic address space pointer to to_private may cause dynamic conversion affecting performance}}
 #endif
 
 
   priv = to_local(glob);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion assigning to 'int *' from 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__private int *__private' from 'int'}}
 #else
-  // expected-error@-4{{assigning '__local int *' to 'int *' changes address space of pointer}}
+  // expected-error@-4{{assigning '__local int *' to '__private int *__private' changes address space of pointer}}
   // expected-warning@-5{{passing non-generic address space pointer to to_local may cause dynamic conversion affecting performance}}
 #endif
 
   glob = to_global(glob);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
   // expected-warning@-4{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}
 #endif
 
   glob = to_private(glob);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
-  // expected-error@-4{{assigning 'int *' to '__global int *' changes address space of pointer}}
+  // expected-error@-4{{assigning '__private int *' to '__global int *__private' changes address space of pointer}}
   // expected-warning@-5{{passing non-generic address space pointer to to_private may cause dynamic conversion affecting performance}}
 #endif
 
   glob = to_local(glob);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
-  // expected-error@-4{{assigning '__local int *' to '__global int *' changes address space of pointer}}
+  // expected-error@-4{{assigning '__local int *' to '__global int *__private' changes address space of pointer}}
   // expected-warning@-5{{passing non-generic address space pointer to to_local may cause dynamic conversion affecting performance}}
 #endif
 
   global char *glob_c = to_global(loc);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion initializing '__global char *' with an expression of type 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion initializing '__global char *__private' with an expression of type 'int'}}
 #else
-  // expected-warning@-4{{incompatible pointer types initializing '__global char *' with an expression of type '__global int *'}}
+  // expected-warning@-4{{incompatible pointer types initializing '__global char *__private' with an expression of type '__global int *'}}
   // expected-warning@-5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}
 #endif
 
   glob_wrong_ty = to_global(glob);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global float *' from 'int'}}
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global float *__private' from 'int'}}
 #else
-  // expected-warning@-4{{incompatible pointer types assigning to '__global float *' from '__global int *'}}
+  // expected-warning@-4{{incompatible pointer types assigning to '__global float *__private' from '__global int *'}}
   // expected-warning@-5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}
 #endif
 
 }

clang/test/SemaOpenCL/vec_step.cl

void foo(int3 arg1, int8 arg2) {
   int res6[vec_step(auto4) == 2 ? 1 : -1];
   int res7[vec_step(int2) == 2 ? 1 : -1];
   int res8[vec_step(int3) == 4 ? 1 : -1];
   int res9[vec_step(int4) == 4 ? 1 : -1];
   int res10[vec_step(int8) == 8 ? 1 : -1];
   int res11[vec_step(int16) == 16 ? 1 : -1];
   int res12[vec_step(void) == 1 ? 1 : -1];
 
-  int res13 = vec_step(*incomplete1); // expected-error {{'vec_step' requires built-in scalar or vector type, 'struct S' invalid}}
-  int res14 = vec_step(int16*); // expected-error {{'vec_step' requires built-in scalar or vector type, 'int16 *' invalid}}
+  int res13 = vec_step(*incomplete1); // expected-error {{'vec_step' requires built-in scalar or vector type, '__private struct S' invalid}}
+  int res14 = vec_step(int16*); // expected-error {{'vec_step' requires built-in scalar or vector type, '__private int16 *' invalid}}
   int res15 = vec_step(void(void)); // expected-error {{'vec_step' requires built-in scalar or vector type, 'void (void)' invalid}}
 }

clang/test/SemaOpenCL/vector_conv_invalid.cl

 // RUN: %clang_cc1 -verify %s
 
 typedef unsigned int uint4 __attribute((ext_vector_type(4)));
 typedef int int4 __attribute((ext_vector_type(4)));
 typedef int int3 __attribute((ext_vector_type(3)));
 typedef unsigned uint3 __attribute((ext_vector_type(3)));
 
 void vector_conv_invalid(const global int4 *const_global_ptr) {
   uint4 u = (uint4)(1);
-  int4 i = u; // expected-error{{initializing 'int4' (vector of 4 'int' values) with an expression of incompatible type 'uint4' (vector of 4 'unsigned int' values)}}
+  int4 i = u; // expected-error{{initializing '__private int4' (vector of 4 'int' values) with an expression of incompatible type '__private uint4' (vector of 4 'unsigned int' values)}}
   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}}
-  global int4 *global_ptr = const_global_ptr;                 // expected-warning {{initializing '__global int4 *' with an expression of type 'const __global int4 *' discards qualifiers}}
+  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 *__private' with an expression of type 'const __global int4 *__private' discards qualifiers}}
   global_ptr = (global int4 *)const_global_ptr;
 }

clang/test/SemaOpenCLCXX/address-space-deduction.cl

 };
 
 
 // Addr spaces for pointee of dependent types are not deduced
 // during parsing but during template instantiation instead.
 
 template <class T>
 struct x1 {
-//CHECK: -CXXMethodDecl {{.*}} operator= 'x1<T> &(const x1<T> &){{( __attribute__.*)?}} __generic'
-//CHECK: -CXXMethodDecl {{.*}} operator= '__generic x1<int> &(const __generic x1<int> &){{( __attribute__.*)?}} __generic'
+//CHECK: -CXXMethodDecl {{.*}} operator= 'x1<T> &(const x1<T> &__private){{( __attribute__.*)?}} __generic'
+//CHECK: -CXXMethodDecl {{.*}} operator= '__generic x1<int> &(const __generic x1<int> &__private){{( __attribute__.*)?}} __generic'
   x1<T>& operator=(const x1<T>& xx) {
     y = xx.y;
     return *this;
   }
   int y;
 };
 
 template <class T>
 struct x2 {
-//CHECK: -CXXMethodDecl {{.*}} foo 'void (x1<T> *){{( __attribute__.*)?}} __generic'
-//CHECK: -CXXMethodDecl {{.*}} foo 'void (__generic x1<int> *){{( __attribute__.*)?}} __generic'
+//CHECK: -CXXMethodDecl {{.*}} foo 'void (x1<T> *__private){{( __attribute__.*)?}} __generic'
+//CHECK: -CXXMethodDecl {{.*}} foo 'void (__generic x1<int> *__private){{( __attribute__.*)?}} __generic'
   void foo(x1<T>* xx) {
     m[0] = *xx;
   }
 //CHECK: -FieldDecl {{.*}}  m 'x1<int> [2]'
   x1<T> m[2];
 };
 
 void bar(__global x1<int> *xx, __global x2<int> *bar) {
   bar->foo(xx);
 }
 
 template <typename T>
 class x3 : public T {
 public:
-  //CHECK: -CXXConstructorDecl {{.*}} x3<T> 'void (const x3<T> &){{( __attribute__.*)?}} __generic'
+  //CHECK: -CXXConstructorDecl {{.*}} x3<T> 'void (const x3<T> &__private){{( __attribute__.*)?}} __generic'
   x3(const x3 &t);
 };
-//CHECK: -CXXConstructorDecl {{.*}} x3<T> 'void (const x3<T> &){{( __attribute__.*)?}} __generic'
+//CHECK: -CXXConstructorDecl {{.*}} x3<T> 'void (const x3<T> &__private){{( __attribute__.*)?}} __generic'
 template <typename T>
 x3<T>::x3(const x3<T> &t) {}
 
 template <class T>
 T xxx(T *in1, T in2) {
   // This pointer can't be deduced to generic because addr space
   // will be taken from the template argument.
-  //CHECK: `-VarDecl {{.*}} i 'T *' cinit
+  //CHECK: `-VarDecl {{.*}} '__private T *__private' cinit
   T *i = in1;
   T ii;
   __private T *ptr = &ii;
   ptr = &in2;
   return *i;
 }
 
 __kernel void test() {
   int foo[10];
   xxx<__private int>(&foo[0], foo[0]);
   // FIXME: Template param deduction fails here because
   // temporaries are not in the __private address space.
   // It is probably reasonable to put them in __private
   // considering that stack and function params are
   // implicitly in __private.
   // However, if temporaries are left in default addr
   // space we should at least pretty print the __private
   // addr space. Otherwise diagnostic apprears to be
   // confusing.
   //xxx(&foo[0], foo[0]);
 }
 
 // Addr space for pointer/reference to an array
-//CHECK: FunctionDecl {{.*}} t1 'void (const float (__generic &)[2])'
+//CHECK: FunctionDecl {{.*}} t1 'void (const float (__generic &__private)[2])'
 void t1(const float (&fYZ)[2]);
-//CHECK: FunctionDecl {{.*}} t2 'void (const float (__generic *)[2])'
+//CHECK: FunctionDecl {{.*}} t2 'void (const float (__generic *__private)[2])'
 void t2(const float (*fYZ)[2]);
-//CHECK: FunctionDecl {{.*}} t3 'void (float (((__generic *)))[2])'
+//CHECK: FunctionDecl {{.*}} t3 'void (float (((__generic *__private)))[2])'
 void t3(float(((*fYZ)))[2]);
-//CHECK: FunctionDecl {{.*}} t4 'void (float (((__generic *__generic *)))[2])'
+//CHECK: FunctionDecl {{.*}} t4 'void (float (((__generic *__generic *__private)))[2])'
 void t4(float(((**fYZ)))[2]);
-//CHECK: FunctionDecl {{.*}} t5 'void (float (__generic *(__generic *))[2])'
+//CHECK: FunctionDecl {{.*}} t5 'void (float (__generic *(__generic *__private))[2])'
 void t5(float (*(*fYZ))[2]);
 
 __kernel void k() {
   __local float x[2];
   __local float(*p)[2];
   t1(x);
   t2(&x);
   t3(&x);
   t4(&p);
   t5(&p);
 }

clang/test/SemaOpenCLCXX/address-space-lambda.cl

 //RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify | FileCheck %s
 
-//CHECK: CXXMethodDecl {{.*}} constexpr operator() 'int (int) const __generic'
+//CHECK: CXXMethodDecl {{.*}} constexpr operator() 'int (__private int) const __generic'
 auto glambda = [](auto a) { return a; };
 
 __kernel void test() {
   int i;
 //CHECK: CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic'
   auto  llambda = [&]() {i++;};
   llambda();
   glambda(1);
   // Test lambda with default parameters
 //CHECK: CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic'
   [&] {i++;} ();
   __constant auto err = [&]() {}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__generic'}}
   err();                          //expected-error-re{{no matching function for call to object of type '__constant (lambda at {{.*}})'}}
   // FIXME: There is very limited addr space functionality
   // we can test when taking lambda type from the object.
   // The limitation is due to addr spaces being added to all
   // objects in OpenCL. Once we add metaprogramming utility
   // for removing address spaces from a type we can enhance
   // testing here.
   (*(__constant decltype(llambda) *)nullptr)(); //expected-error{{multiple address spaces specified for type}}
   (*(decltype(llambda) *)nullptr)();
 }
 
 __kernel void test_qual() {
-//CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const'
+//CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const __private'
   auto priv1 = []() __private {};
   priv1();
 //CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic'
   auto priv2 = []() __generic {};
   priv2();
-  auto priv3 = []() __global {}; //expected-note{{candidate function not viable: 'this' object is in default address space, but method expects object in address space '__global'}} //expected-note{{conversion candidate of type 'void (*)()'}}
+  auto priv3 = []() __global {}; //expected-note{{candidate function not viable: 'this' object is in address space '__private', but method expects object in address space '__global'}} //expected-note{{conversion candidate of type 'void (*)()'}}
   priv3(); //expected-error{{no matching function for call to object of type}}
 
-  __constant auto const1 = []() __private{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in default address space}} //expected-note{{conversion candidate of type 'void (*)()'}}
+  __constant auto const1 = []() __private{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__private'}} //expected-note{{conversion candidate of type 'void (*)()'}}
   const1(); //expected-error{{no matching function for call to object of type '__constant (lambda at}}
   __constant auto const2 = []() __generic{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__generic'}} //expected-note{{conversion candidate of type 'void (*)()'}}
   const2(); //expected-error{{no matching function for call to object of type '__constant (lambda at}}
 //CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const __constant'
   __constant auto const3 = []() __constant{};
   const3();
 
   [&] () __global {} (); //expected-error{{no matching function for call to object of type '(lambda at}} expected-note{{candidate function not viable: 'this' object is in default address space, but method expects object in address space '__global'}}
-  [&] () __private {} (); //expected-error{{no matching function for call to object of type '(lambda at}} expected-note{{candidate function not viable: 'this' object is in default address space, but method expects object in default address space}}
+  [&] () __private {} (); //expected-error{{no matching function for call to object of type '(lambda at}} expected-note{{candidate function not viable: 'this' object is in default address space, but method expects object in address space '__private'}}
 
   [&] __private {} (); //expected-error{{lambda requires '()' before attribute specifier}} expected-error{{expected body of lambda expression}}
 
   [&] () mutable __private {} ();
   [&] () __private mutable {} (); //expected-error{{expected body of lambda expression}}
 }
 

clang/test/SemaOpenCLCXX/address-space-templates.cl

 
 template <typename T>
 T *foo2(T *i) {
   return i;
 }
 
 template <typename T>
 void foo3() {
-  __private T ii; // expected-error{{conflicting address space qualifiers are provided between types 'T' and '__global int'}}
+  __private T ii; // expected-error{{conflicting address space qualifiers are provided between types '__private T' and '__global int'}}
 }
 
 void bar() {
   S<const __global int> sintgl; // expected-note{{in instantiation of template class 'S<const __global int>' requested here}}
 
   foo1<__local int>(1); // expected-error{{no matching function for call to 'foo1'}}
   foo2<__global int>(0);
   foo3<__global int>(); // expected-note{{in instantiation of function template specialization 'foo3<__global int>' requested here}}
 }

clang/test/SemaOpenCLCXX/addrspace-auto.cl

 //RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify | FileCheck %s
 
 __constant int i = 1;
 //CHECK: |-VarDecl {{.*}} ai '__global int':'__global int'
 auto ai = i;
 
 kernel void test() {
   int i;
-  //CHECK: VarDecl {{.*}} ai 'int':'int'
+  //CHECK: VarDecl {{.*}} ai '__private int':'__private int'
   auto ai = i;
 
   constexpr int c = 1;
   //CHECK: VarDecl {{.*}} used cai '__constant int':'__constant int'
   __constant auto cai = c;
-  //CHECK: VarDecl {{.*}} aii 'int':'int'
+  //CHECK: VarDecl {{.*}} aii '__private int':'__private int'
   auto aii = cai;
 
-  //CHECK: VarDecl {{.*}} ref 'int &'
+  //CHECK: VarDecl {{.*}} ref '__private int &__private'
   auto &ref = i;
-  //CHECK: VarDecl {{.*}} ptr 'int *'
+  //CHECK: VarDecl {{.*}} ptr '__private int *__private'
   auto *ptr = &i;
-  //CHECK: VarDecl {{.*}} ref_c '__constant int &'
+  //CHECK: VarDecl {{.*}} ref_c '__constant int &__private'
   auto &ref_c = cai;
 
-  //CHECK: VarDecl {{.*}} ptrptr 'int *__generic *'
+  //CHECK: VarDecl {{.*}} ptrptr '__private int *__generic *__private'
   auto **ptrptr = &ptr;
-  //CHECK: VarDecl {{.*}} refptr 'int *__generic &'
+  //CHECK: VarDecl {{.*}} refptr '__private int *__generic &__private'
   auto *&refptr = ptr;
 
-  //CHECK: VarDecl {{.*}} invalid gref '__global auto &'
-  __global auto &gref = i; //expected-error{{variable 'gref' with type '__global auto &' has incompatible initializer of type 'int'}}
+  //CHECK: VarDecl {{.*}} invalid gref '__global auto &__private'
+  __global auto &gref = i; //expected-error{{variable 'gref' with type '__global auto &__private' has incompatible initializer of type '__private int'}}
   __local int *ptr_l;
-  //CHECK: VarDecl {{.*}} invalid gptr '__global auto *'
-  __global auto *gptr = ptr_l; //expected-error{{variable 'gptr' with type '__global auto *' has incompatible initializer of type '__local int *'}}
+  //CHECK: VarDecl {{.*}} invalid gptr '__global auto *__private'
+  __global auto *gptr = ptr_l; //expected-error{{variable 'gptr' with type '__global auto *__private' has incompatible initializer of type '__local int *__private'}}
 }