[clang] 869d17d - [OpenCL] Pretty print __private addr space

[Anastasia Stulova via cfe-commits]

Author: Anastasia Stulova
Date: 2019-12-27T13:42:07Z
New Revision: 869d17d851ba08fc35ade3d1077b4986de401e89

URL: https://github.com/llvm/llvm-project/commit/869d17d851ba08fc35ade3d1077b4986de401e89
DIFF: https://github.com/llvm/llvm-project/commit/869d17d851ba08fc35ade3d1077b4986de401e89.diff

LOG: [OpenCL] Pretty print __private addr space

Add printing of __private address space to TypePrinter to allow
it appears in diagnostics and AST dumps as all other language
addr spaces.

Tags: #clang

Differential Revision: https://reviews.llvm.org/D71272

Added: 
    

Modified: 
    clang/lib/AST/TypePrinter.cpp
    clang/lib/Sema/SemaDecl.cpp
    clang/test/AST/language_address_space_attribute.cpp
    clang/test/Index/opencl-types.cl
    clang/test/Parser/opencl-astype.cl
    clang/test/Parser/opencl-atomics-cl20.cl
    clang/test/SemaOpenCL/access-qualifier.cl
    clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
    clang/test/SemaOpenCL/address-spaces.cl
    clang/test/SemaOpenCL/arithmetic-conversions.cl
    clang/test/SemaOpenCL/as_type.cl
    clang/test/SemaOpenCL/atomic-ops.cl
    clang/test/SemaOpenCL/cl20-device-side-enqueue.cl
    clang/test/SemaOpenCL/clk_event_t.cl
    clang/test/SemaOpenCL/event_t.cl
    clang/test/SemaOpenCL/extension-begin.cl
    clang/test/SemaOpenCL/half.cl
    clang/test/SemaOpenCL/images.cl
    clang/test/SemaOpenCL/intel-subgroup-avc-ext-types.cl
    clang/test/SemaOpenCL/invalid-block.cl
    clang/test/SemaOpenCL/invalid-image.cl
    clang/test/SemaOpenCL/invalid-kernel-parameters.cl
    clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl
    clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl
    clang/test/SemaOpenCL/null_literal.cl
    clang/test/SemaOpenCL/null_queue.cl
    clang/test/SemaOpenCL/predefined-expr.cl
    clang/test/SemaOpenCL/queue_t_overload.cl
    clang/test/SemaOpenCL/shifts.cl
    clang/test/SemaOpenCL/to_addr_builtin.cl
    clang/test/SemaOpenCL/vec_step.cl
    clang/test/SemaOpenCL/vector_conv_invalid.cl
    clang/test/SemaOpenCLCXX/address-space-deduction.cl
    clang/test/SemaOpenCLCXX/address-space-lambda.cl
    clang/test/SemaOpenCLCXX/address-space-templates.cl
    clang/test/SemaOpenCLCXX/addrspace-auto.cl

Removed: 
    


################################################################################
diff  --git a/clang/lib/AST/TypePrinter.cpp b/clang/lib/AST/TypePrinter.cpp
index be23a497577e..c2f4baec989e 100644
--- a/clang/lib/AST/TypePrinter.cpp
+++ b/clang/lib/AST/TypePrinter.cpp
@@ -1784,7 +1784,7 @@ std::string Qualifiers::getAddrSpaceAsString(LangAS AS) {
   case LangAS::opencl_local:
     return "__local";
   case LangAS::opencl_private:
-    return "";
+    return "__private";
   case LangAS::opencl_constant:
     return "__constant";
   case LangAS::opencl_generic:

diff  --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 8f68be716bd3..9ca7a80d3ada 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -8381,7 +8381,7 @@ static bool isOpenCLSizeDependentType(ASTContext &C, QualType Ty) {
   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;
 

diff  --git a/clang/test/AST/language_address_space_attribute.cpp b/clang/test/AST/language_address_space_attribute.cpp
index 7c6bdca06c06..f9ab4612320b 100644
--- a/clang/test/AST/language_address_space_attribute.cpp
+++ b/clang/test/AST/language_address_space_attribute.cpp
@@ -22,10 +22,10 @@ void langas() {
   // 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 *'

diff  --git a/clang/test/Index/opencl-types.cl b/clang/test/Index/opencl-types.cl
index e132c9d05ac8..496f38752fa2 100644
--- a/clang/test/Index/opencl-types.cl
+++ b/clang/test/Index/opencl-types.cl
@@ -16,12 +16,12 @@ void kernel testFloatTypes() {
   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
 
@@ -38,18 +38,18 @@ void kernel OCLImage2dMSAADepthROTest(read_only image2d_msaa_depth_t scalarOCLIm
 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);
@@ -64,18 +64,18 @@ void kernel OCLImage2dMSAADepthWOTest(write_only image2d_msaa_depth_t scalarOCLI
 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);
@@ -90,24 +90,24 @@ void kernel OCLImage2dMSAADepthRWTest(read_write image2d_msaa_depth_t scalarOCLI
 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
@@ -121,9 +121,9 @@ void kernel testMiscOpenCLTypes() {
 }
 
 // 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
 
@@ -131,4 +131,4 @@ 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]

diff  --git a/clang/test/Parser/opencl-astype.cl b/clang/test/Parser/opencl-astype.cl
index 903c42ee8c14..4a26fb635ddf 100644
--- a/clang/test/Parser/opencl-astype.cl
+++ b/clang/test/Parser/opencl-astype.cl
@@ -11,7 +11,7 @@ void test_astype() {
   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;

diff  --git a/clang/test/Parser/opencl-atomics-cl20.cl b/clang/test/Parser/opencl-atomics-cl20.cl
index ad67db0bab8a..844902e847f7 100644
--- a/clang/test/Parser/opencl-atomics-cl20.cl
+++ b/clang/test/Parser/opencl-atomics-cl20.cl
@@ -66,9 +66,9 @@ 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

diff  --git a/clang/test/SemaOpenCL/access-qualifier.cl b/clang/test/SemaOpenCL/access-qualifier.cl
index 1753a1dbefdd..a5e1b65daf70 100644
--- a/clang/test/SemaOpenCL/access-qualifier.cl
+++ b/clang/test/SemaOpenCL/access-qualifier.cl
@@ -25,11 +25,11 @@ void myReadWrite(read_write image1d_t); // expected-error {{access qualifier 're
 
 
 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) {
@@ -43,7 +43,7 @@ kernel void k4(img1d_rw 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) {
@@ -71,7 +71,7 @@ kernel void k13(__read_write image1d_t i){} // expected-error{{access qualifier
 #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
 
@@ -93,7 +93,7 @@ kernel void pipe_ro_twice_typedef(read_only ROPipeInt i){} // expected-warning{{
 // expected-note at -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 at -25 {{passing argument to parameter here}}
 }
 #endif

diff  --git a/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl b/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
index 0209e1ea24de..14f7cf3a0e7e 100644
--- a/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
+++ b/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
@@ -66,7 +66,7 @@ void f_priv(__private int *arg_priv) {}
 #if !__OPENCL_CPP_VERSION__
 // expected-note at -2{{passing argument to parameter 'arg_priv' here}}
 #else
-// expected-note-re at -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 at -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) {}
@@ -85,45 +85,45 @@ void test_conversion(__global int *arg_glob, __local int *arg_loc,
   AS int *var_init1 = arg_glob;
 #ifdef CONSTANT
 #if !__OPENCL_CPP_VERSION__
-// expected-error at -3{{initializing '__constant int *' with an expression of type '__global int *' changes address space of pointer}}
+// expected-error at -3{{initializing '__constant int *__private' with an expression of type '__global int *__private' changes address space of pointer}}
 #else
-// expected-error at -5{{cannot initialize a variable of type '__constant int *' with an lvalue of type '__global int *'}}
+// expected-error at -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 at -3{{initializing '__{{global|constant}} int *' with an expression of type '__local int *' changes address space of pointer}}
+// expected-error-re at -3{{initializing '__{{global|constant}} int *__private' with an expression of type '__local int *__private' changes address space of pointer}}
 #else
-// expected-error-re at -5{{cannot initialize a variable of type '__{{global|constant}} int *' with an lvalue of type '__local int *'}}
+// expected-error-re at -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 at -3{{initializing '__{{global|generic}} int *' with an expression of type '__constant int *' changes address space of pointer}}
+// expected-error-re at -3{{initializing '__{{global|generic}} int *__private' with an expression of type '__constant int *__private' changes address space of pointer}}
 #else
-// expected-error-re at -5{{cannot initialize a variable of type '__{{global|generic}} int *' with an lvalue of type '__constant int *'}}
+// expected-error-re at -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 at -3{{initializing '__{{global|constant}} int *' with an expression of type 'int *' changes address space of pointer}}
+// expected-error-re at -3{{initializing '__{{global|constant}} int *__private' with an expression of type '__private int *__private' changes address space of pointer}}
 #else
-// expected-error-re at -5{{cannot initialize a variable of type '__{{global|constant}} int *' with an lvalue of type 'int *'}}
+// expected-error-re at -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 at -3{{initializing '__{{global|constant}} int *' with an expression of type '__generic int *' changes address space of pointer}}
+// expected-error-re at -3{{initializing '__{{global|constant}} int *__private' with an expression of type '__generic int *__private' changes address space of pointer}}
 #else
-// expected-error-re at -5{{cannot initialize a variable of type '__{{global|constant}} int *' with an lvalue of type '__generic int *'}}
+// expected-error-re at -5{{cannot initialize a variable of type '__{{global|constant}} int *__private' with an lvalue of type '__generic int *__private'}}
 #endif
 #endif
 
@@ -157,9 +157,9 @@ void test_conversion(__global int *arg_glob, __local int *arg_loc,
   AS int *var_cast4 = (AS int *)arg_priv;
 #ifndef GENERIC
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re at -3{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re at -3{{casting '__private int *' to type '__{{global|constant}} int *' changes address space of pointer}}
 #else
-// expected-error-re at -5{{C-style cast from 'int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
+// expected-error-re at -5{{C-style cast from '__private int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
 #endif
 #endif
 
@@ -176,45 +176,45 @@ void test_conversion(__global int *arg_glob, __local int *arg_loc,
   var_impl = arg_glob;
 #ifdef CONSTANT
 #if !__OPENCL_CPP_VERSION__
-// expected-error at -3{{assigning '__global int *' to '__constant int *' changes address space of pointer}}
+// expected-error at -3{{assigning '__global int *__private' to '__constant int *__private' changes address space of pointer}}
 #else
-// expected-error at -5{{assigning to '__constant int *' from incompatible type '__global int *'}}
+// expected-error at -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 at -3{{assigning '__local int *' to '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re at -3{{assigning '__local int *__private' to '__{{global|constant}} int *__private' changes address space of pointer}}
 #else
-// expected-error-re at -5{{assigning to '__{{global|constant}} int *' from incompatible type '__local int *'}}
+// expected-error-re at -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 at -3{{assigning '__constant int *' to '__{{global|generic}} int *' changes address space of pointer}}
+// expected-error-re at -3{{assigning '__constant int *__private' to '__{{global|generic}} int *__private' changes address space of pointer}}
 #else
-// expected-error-re at -5{{assigning to '__{{global|generic}} int *' from incompatible type '__constant int *'}}
+// expected-error-re at -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 at -3{{assigning 'int *' to '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re at -3{{assigning '__private int *__private' to '__{{global|constant}} int *__private' changes address space of pointer}}
 #else
-// expected-error-re at -5{{assigning to '__{{global|constant}} int *' from incompatible type 'int *'}}
+// expected-error-re at -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 at -3{{assigning '__generic int *' to '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re at -3{{assigning '__generic int *__private' to '__{{global|constant}} int *__private' changes address space of pointer}}
 #else
-// expected-error-re at -5{{assigning to '__{{global|constant}} int *' from incompatible type '__generic int *'}}
+// expected-error-re at -5{{assigning to '__{{global|constant}} int *' from incompatible type '__generic int *__private'}}
 #endif
 #endif
 
@@ -248,9 +248,9 @@ void test_conversion(__global int *arg_glob, __local int *arg_loc,
   var_cast4 = (AS int *)arg_priv;
 #ifndef GENERIC
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re at -3{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re at -3{{casting '__private int *' to type '__{{global|constant}} int *' changes address space of pointer}}
 #else
-// expected-error-re at -5{{C-style cast from 'int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
+// expected-error-re at -5{{C-style cast from '__private int *' to '__{{global|constant}} int *' converts between mismatching address spaces}}
 #endif
 #endif
 
@@ -294,9 +294,9 @@ void test_conversion(__global int *arg_glob, __local int *arg_loc,
   b = var_cmp <= arg_priv;
 #ifndef GENERIC
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re at -3{{comparison between  ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+// expected-error-re at -3{{comparison between  ('__{{global|constant}} int *' and '__private int *') which are pointers to non-overlapping address spaces}}
 #else
-// expected-error-re at -5{{comparison of distinct pointer types ('__{{global|constant}} int *' and 'int *')}}
+// expected-error-re at -5{{comparison of distinct pointer types ('__{{global|constant}} int *' and '__private int *')}}
 #endif
 #endif
 
@@ -327,7 +327,7 @@ void test_conversion(__global int *arg_glob, __local int *arg_loc,
 
   b = var_sub - arg_priv;
 #ifndef GENERIC
-// expected-error-re at -2{{arithmetic operation with operands of type  ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+// expected-error-re at -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;
@@ -338,7 +338,7 @@ void test_conversion(__global int *arg_glob, __local int *arg_loc,
   f_glob(var_sub);
 #ifndef GLOBAL
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re at -3{{passing '__{{constant|generic}} int *' to parameter of type '__global int *' changes address space of pointer}}
+// expected-error-re at -3{{passing '__{{constant|generic}} int *__private' to parameter of type '__global int *' changes address space of pointer}}
 #else
 // expected-error at -5{{no matching function for call to 'f_glob'}}
 #endif
@@ -346,7 +346,7 @@ void test_conversion(__global int *arg_glob, __local int *arg_loc,
 
   f_loc(var_sub);
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re at -2{{passing '__{{global|constant|generic}} int *' to parameter of type '__local int *' changes address space of pointer}}
+// expected-error-re at -2{{passing '__{{global|constant|generic}} int *__private' to parameter of type '__local int *' changes address space of pointer}}
 #else
 // expected-error at -4{{no matching function for call to 'f_loc'}}
 #endif
@@ -354,7 +354,7 @@ void test_conversion(__global int *arg_glob, __local int *arg_loc,
   f_const(var_sub);
 #ifndef CONSTANT
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re at -3{{passing '__{{global|generic}} int *' to parameter of type '__constant int *' changes address space of pointer}}
+// expected-error-re at -3{{passing '__{{global|generic}} int *__private' to parameter of type '__constant int *' changes address space of pointer}}
 #else
 // expected-error at -5{{no matching function for call to 'f_const'}}
 #endif
@@ -362,7 +362,7 @@ void test_conversion(__global int *arg_glob, __local int *arg_loc,
 
   f_priv(var_sub);
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re at -2{{passing '__{{global|constant|generic}} int *' to parameter of type 'int *' changes address space of pointer}}
+// expected-error-re at -2{{passing '__{{global|constant|generic}} int *__private' to parameter of type '__private int *' changes address space of pointer}}
 #else
 // expected-error at -4{{no matching function for call to 'f_priv'}}
 #endif
@@ -370,7 +370,7 @@ void test_conversion(__global int *arg_glob, __local int *arg_loc,
   f_gen(var_sub);
 #ifdef CONSTANT
 #if !__OPENCL_CPP_VERSION__
-// expected-error at -3{{passing '__constant int *' to parameter of type '__generic int *' changes address space of pointer}}
+// expected-error at -3{{passing '__constant int *__private' to parameter of type '__generic int *' changes address space of pointer}}
 #else
 // expected-error at -5{{no matching function for call to 'f_gen'}}
 #endif
@@ -414,9 +414,9 @@ void test_ternary() {
   var_gen = 0 ? var_cond : var_priv;
 #ifndef GENERIC
 #if !__OPENCL_CPP_VERSION__
-// expected-error-re at -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 at -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 at -5{{incompatible operand types ('__{{global|constant}} int *' and 'int *')}}
+// expected-error-re at -5{{incompatible operand types ('__{{global|constant}} int *' and '__private int *')}}
 #endif
 #endif
 
@@ -470,12 +470,12 @@ void test_ternary() {
   __private char *var_priv_ch;
   var_void_gen = 0 ? var_cond : var_priv_ch;
 #if __OPENCL_CPP_VERSION__
-// expected-error-re at -2{{incompatible operand types ('__{{constant|global|generic}} int *' and 'char *')}}
+// expected-error-re at -2{{incompatible operand types ('__{{constant|global|generic}} int *' and '__private char *')}}
 #else
 #ifndef GENERIC
-// expected-error-re at -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 at -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 at -7{{pointer type mismatch ('__generic int *' and 'char *')}}
+// expected-warning at -7{{pointer type mismatch ('__generic int *' and '__private char *')}}
 #endif
 #endif
 
@@ -517,7 +517,7 @@ void test_pointer_chains() {
   var_as_as_int = var_asc_asc_int;
 #if !__OPENCL_CPP_VERSION__
 #ifdef GENERIC
-// expected-error at -3 {{assigning '__local int *__local *' to '__generic int *__generic *' changes address space of nested pointer}}
+// expected-error at -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;

diff  --git a/clang/test/SemaOpenCL/address-spaces.cl b/clang/test/SemaOpenCL/address-spaces.cl
index a28069470177..bc9a07d88704 100644
--- a/clang/test/SemaOpenCL/address-spaces.cl
+++ b/clang/test/SemaOpenCL/address-spaces.cl
@@ -10,15 +10,15 @@ __kernel void foo(__global int *gip) {
 
   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 at -2 {{assigning '__constant int *' to '__generic int *' changes address space of pointer}}
+// expected-error at -2 {{assigning '__constant int *' to '__generic int *__private' changes address space of pointer}}
 #else
 // expected-error at -4 {{assigning to '__generic int *' from incompatible type '__constant int *'}}
 #endif
@@ -46,9 +46,9 @@ void explicit_cast(__global int *g, __local int *l, __constant int *c, __private
 #endif
   g = (__global int *)p;
 #if !__OPENCL_CPP_VERSION__
-// expected-error at -2 {{casting 'int *' to type '__global int *' changes address space of pointer}}
+// expected-error at -2 {{casting '__private int *' to type '__global int *' changes address space of pointer}}
 #else
-// expected-error at -4 {{C-style cast from 'int *' to '__global int *' converts between mismatching address spaces}}
+// expected-error at -4 {{C-style cast from '__private int *' to '__global int *' converts between mismatching address spaces}}
 #endif
   l = (__local int *)g;
 #if !__OPENCL_CPP_VERSION__
@@ -70,9 +70,9 @@ void explicit_cast(__global int *g, __local int *l, __constant int *c, __private
 #endif
   l = (__local int *)p;
 #if !__OPENCL_CPP_VERSION__
-// expected-error at -2 {{casting 'int *' to type '__local int *' changes address space of pointer}}
+// expected-error at -2 {{casting '__private int *' to type '__local int *' changes address space of pointer}}
 #else
-// expected-error at -4 {{C-style cast from 'int *' to '__local int *' converts between mismatching address spaces}}
+// expected-error at -4 {{C-style cast from '__private int *' to '__local int *' converts between mismatching address spaces}}
 #endif
   c = (__constant int *)g;
 #if !__OPENCL_CPP_VERSION__
@@ -88,33 +88,33 @@ void explicit_cast(__global int *g, __local int *l, __constant int *c, __private
 #endif
   c = (__constant int *)p;
 #if !__OPENCL_CPP_VERSION__
-// expected-error at -2 {{casting 'int *' to type '__constant int *' changes address space of pointer}}
+// expected-error at -2 {{casting '__private int *' to type '__constant int *' changes address space of pointer}}
 #else
-// expected-error at -4 {{C-style cast from 'int *' to '__constant int *' converts between mismatching address spaces}}
+// expected-error at -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 at -2 {{casting '__global int *' to type 'int *' changes address space of pointer}}
+// expected-error at -2 {{casting '__global int *' to type '__private int *' changes address space of pointer}}
 #else
-// expected-error at -4 {{C-style cast from '__global int *' to 'int *' converts between mismatching address spaces}}
+// expected-error at -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 at -2 {{casting '__local int *' to type 'int *' changes address space of pointer}}
+// expected-error at -2 {{casting '__local int *' to type '__private int *' changes address space of pointer}}
 #else
-// expected-error at -4 {{C-style cast from '__local int *' to 'int *' converts between mismatching address spaces}}
+// expected-error at -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 at -2 {{casting '__constant int *' to type 'int *' changes address space of pointer}}
+// expected-error at -2 {{casting '__constant int *' to type '__private int *' changes address space of pointer}}
 #else
-// expected-error at -4 {{C-style cast from '__constant int *' to 'int *' converts between mismatching address spaces}}
+// expected-error at -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 at -2 {{casting 'const __constant int *' to type 'int *' changes address space of pointer}}
+// expected-error at -2 {{casting 'const __constant int *' to type '__private int *' changes address space of pointer}}
 #else
-// expected-error at -4 {{C-style cast from 'const __constant int *' to 'int *' converts between mismatching address spaces}}
+// expected-error at -4 {{C-style cast from 'const __constant int *' to '__private int *' converts between mismatching address spaces}}
 #endif
 }
 
@@ -126,98 +126,98 @@ void ok_explicit_casts(__global int *g, __global int *g2, __local int *l, __loca
 
 #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}}
+  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 *' 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 **'}}
+  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}}
+  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 *' 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}}
+  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 *' 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}}
+  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}}
+  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 **' 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 **'}}
+  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}}
+  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 **' 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 **'}}
+  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;

diff  --git a/clang/test/SemaOpenCL/arithmetic-conversions.cl b/clang/test/SemaOpenCL/arithmetic-conversions.cl
index 23103caf85ef..3d9f34851e90 100644
--- a/clang/test/SemaOpenCL/arithmetic-conversions.cl
+++ b/clang/test/SemaOpenCL/arithmetic-conversions.cl
@@ -14,10 +14,10 @@ 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))}}
 }

diff  --git a/clang/test/SemaOpenCL/as_type.cl b/clang/test/SemaOpenCL/as_type.cl
index ad418097d902..e43ec7eb5f06 100644
--- a/clang/test/SemaOpenCL/as_type.cl
+++ b/clang/test/SemaOpenCL/as_type.cl
@@ -1,14 +1,14 @@
 // 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}}
 }

diff  --git a/clang/test/SemaOpenCL/atomic-ops.cl b/clang/test/SemaOpenCL/atomic-ops.cl
index d8021e33af65..c95c73d52e52 100644
--- a/clang/test/SemaOpenCL/atomic-ops.cl
+++ b/clang/test/SemaOpenCL/atomic-ops.cl
@@ -82,13 +82,13 @@ void f(atomic_int *i, const atomic_int *ci,
 
   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 
diff erent 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);
@@ -107,7 +107,7 @@ void memory_checks(atomic_int *Ap, int *p, int val) {
   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);
@@ -187,7 +187,7 @@ void synchscope_checks(atomic_int *Ap, int scope) {
   (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) {

diff  --git a/clang/test/SemaOpenCL/cl20-device-side-enqueue.cl b/clang/test/SemaOpenCL/cl20-device-side-enqueue.cl
index 7e8ec3d498cc..f63e2913c749 100644
--- a/clang/test/SemaOpenCL/cl20-device-side-enqueue.cl
+++ b/clang/test/SemaOpenCL/cl20-device-side-enqueue.cl
@@ -91,7 +91,7 @@ kernel void enqueue_kernel_tests() {
                  },
                  c, 1024L);
 #ifdef WCONV
-// expected-warning-re at -2{{implicit conversion changes signedness: 'char' to 'unsigned {{int|long}}'}}
+// expected-warning-re at -2{{implicit conversion changes signedness: '__private char' to 'unsigned {{int|long}}'}}
 #endif
 #define UINT_MAX 4294967295
 

diff  --git a/clang/test/SemaOpenCL/clk_event_t.cl b/clang/test/SemaOpenCL/clk_event_t.cl
index 14a9318fa45c..2259359ae3c9 100644
--- a/clang/test/SemaOpenCL/clk_event_t.cl
+++ b/clang/test/SemaOpenCL/clk_event_t.cl
@@ -12,7 +12,7 @@ int clk_event_tests() {
   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;
   }
 

diff  --git a/clang/test/SemaOpenCL/event_t.cl b/clang/test/SemaOpenCL/event_t.cl
index ab7f09170e9c..80af4013a6cb 100644
--- a/clang/test/SemaOpenCL/event_t.cl
+++ b/clang/test/SemaOpenCL/event_t.cl
@@ -1,6 +1,6 @@
 // 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}}
@@ -8,7 +8,7 @@ constant struct evt_s {
 
 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);

diff  --git a/clang/test/SemaOpenCL/extension-begin.cl b/clang/test/SemaOpenCL/extension-begin.cl
index 367652ccdd22..fdb481f8e092 100644
--- a/clang/test/SemaOpenCL/extension-begin.cl
+++ b/clang/test/SemaOpenCL/extension-begin.cl
@@ -44,7 +44,7 @@ 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 at extension-begin.h:18 {{candidate unavailable as it requires OpenCL extension 'my_ext' to be enabled}}

diff  --git a/clang/test/SemaOpenCL/half.cl b/clang/test/SemaOpenCL/half.cl
index e8da3a2a2137..2ea971c08510 100644
--- a/clang/test/SemaOpenCL/half.cl
+++ b/clang/test/SemaOpenCL/half.cl
@@ -4,13 +4,13 @@
 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;
@@ -26,7 +26,7 @@ half half_disabled(half *p, // expected-error{{declaring function return value o
 }
 
 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

diff  --git a/clang/test/SemaOpenCL/images.cl b/clang/test/SemaOpenCL/images.cl
index 9d6092c8c916..d9305dd389b6 100644
--- a/clang/test/SemaOpenCL/images.cl
+++ b/clang/test/SemaOpenCL/images.cl
@@ -7,26 +7,26 @@ void img2d_default(image2d_t); // expected-note 2{{passing argument to parameter
 
 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'}}
 }

diff  --git a/clang/test/SemaOpenCL/intel-subgroup-avc-ext-types.cl b/clang/test/SemaOpenCL/intel-subgroup-avc-ext-types.cl
index 3392e80565f1..e76d54763016 100644
--- a/clang/test/SemaOpenCL/intel-subgroup-avc-ext-types.cl
+++ b/clang/test/SemaOpenCL/intel-subgroup-avc-ext-types.cl
@@ -16,34 +16,34 @@ 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 at -1 {{initializing 'intel_sub_group_avc_mce_payload_t' with an expression of incompatible type 'int'}}
+  // expected-error at -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 at -1 {{initializing 'intel_sub_group_avc_ime_payload_t' with an expression of incompatible type 'int'}}
+  // expected-error at -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 at -1 {{initializing 'intel_sub_group_avc_ref_payload_t' with an expression of incompatible type 'float'}}
+  // expected-error at -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 at -1 {{initializing 'intel_sub_group_avc_sic_payload_t' with an expression of incompatible type 'struct st'}}
+  // expected-error at -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 at -1 {{initializing 'intel_sub_group_avc_mce_result_t' with an expression of incompatible type 'int'}}
+  // expected-error at -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 at -1 {{initializing 'intel_sub_group_avc_ime_result_t' with an expression of incompatible type 'int'}}
+  // expected-error at -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 at -1 {{initializing 'intel_sub_group_avc_ref_result_t' with an expression of incompatible type 'float'}}
+  // expected-error at -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 at -1 {{initializing 'intel_sub_group_avc_sic_result_t' with an expression of incompatible type 'struct st'}}
+  // expected-error at -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 at -1 {{initializing 'intel_sub_group_avc_ime_result_single_reference_streamout_t' with an expression of incompatible type 'void *'}}
+  // expected-error at -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 at -1 {{excess elements in struct initializer}}
   intel_sub_group_avc_ime_dual_reference_streamin_t dstreamin_list2 = {};
   // expected-error at -1 {{scalar initializer cannot be empty}}
   intel_sub_group_avc_ime_single_reference_streamin_t dstreamin_list3 = {c};
-  // expected-error at -1 {{initializing 'intel_sub_group_avc_ime_single_reference_streamin_t' with an expression of incompatible type 'char'}}
+  // expected-error at -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 at -1 {{initializing 'intel_sub_group_avc_ime_dual_reference_streamin_t' with an expression of incompatible type 'int'}}
+  // expected-error at -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
@@ -53,12 +53,12 @@ void far() {
 
   intel_sub_group_avc_ime_payload_t payload_ime;
   intel_sub_group_avc_ref_payload_t payload_ref = payload_ime;
-  // expected-error at -1 {{initializing 'intel_sub_group_avc_ref_payload_t' with an expression of incompatible type 'intel_sub_group_avc_ime_payload_t'}}
+  // expected-error at -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 at -1 {{assigning to 'intel_sub_group_avc_sic_result_t' from incompatible type 'intel_sub_group_avc_ime_result_t'}}
+  // expected-error at -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

diff  --git a/clang/test/SemaOpenCL/invalid-block.cl b/clang/test/SemaOpenCL/invalid-block.cl
index 7cbcea96d0ac..ec74d16cc9b8 100644
--- a/clang/test/SemaOpenCL/invalid-block.cl
+++ b/clang/test/SemaOpenCL/invalid-block.cl
@@ -12,7 +12,7 @@ void f1() {
   };
   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}}
 }
 
@@ -53,18 +53,18 @@ void f5(int i) {
   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() {

diff  --git a/clang/test/SemaOpenCL/invalid-image.cl b/clang/test/SemaOpenCL/invalid-image.cl
index 9185cc371ebc..a996a67fdfa8 100644
--- a/clang/test/SemaOpenCL/invalid-image.cl
+++ b/clang/test/SemaOpenCL/invalid-image.cl
@@ -5,12 +5,12 @@
 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}}

diff  --git a/clang/test/SemaOpenCL/invalid-kernel-parameters.cl b/clang/test/SemaOpenCL/invalid-kernel-parameters.cl
index e3372b188877..48de39d0f87e 100644
--- a/clang/test/SemaOpenCL/invalid-kernel-parameters.cl
+++ b/clang/test/SemaOpenCL/invalid-kernel-parameters.cl
@@ -1,6 +1,6 @@
 // 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
 
@@ -9,37 +9,37 @@ kernel void half_arg(half x) { } // expected-error{{declaring function parameter
 // bool, half, size_t, ptr
diff _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 at -1{{'size_t' (aka 'unsigned int') declared here}}
-typedef __PTRDIFF_TYPE__ ptr
diff _t; // expected-note{{'ptr
diff _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__ ptr
diff _t; // expected-note{{'__private ptr
diff _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 ptr
diff _t_arg(ptr
diff _t x) {} // expected-error{{'ptr
diff _t' (aka 'int') cannot be used as the type of a kernel parameter}}
+kernel void ptr
diff _t_arg(ptr
diff _t x) {} // expected-error{{'__private ptr
diff _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) { }
@@ -49,7 +49,7 @@ typedef struct ContainsBool // expected-note{{within field of type 'ContainsBool
   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}}
 
 
 
@@ -65,14 +65,14 @@ kernel void image_in_struct_arg(FooImage2D arg) { } // expected-error{{struct ke
 
 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}}
@@ -82,7 +82,7 @@ typedef struct NestedPointer // expected-note 2 {{within field of type 'NestedPo
   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;
 
@@ -96,7 +96,7 @@ struct NestedPointerComplex // expected-note{{within field of type 'NestedPointe
   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;
@@ -114,10 +114,10 @@ typedef struct NestedBool // expected-note 2 {{within field of type 'NestedBool'
   } 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
@@ -132,7 +132,7 @@ typedef struct NestedBool2 // expected-note{{within field of type 'NestedBool2'
   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
@@ -167,8 +167,8 @@ kernel void pointer_in_nested_struct_arg_2(struct Valid valid, struct NestedPoin
 
 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 at -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 at -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}}
 

diff  --git a/clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl b/clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl
index 619b359c7af9..36e76621d24a 100644
--- a/clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl
+++ b/clang/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl
@@ -12,22 +12,22 @@ void test1(read_only pipe int p, global int* 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)}}
 }
 
@@ -41,22 +41,22 @@ void test2(write_only pipe int p, global int* 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)}}
 }
 

diff  --git a/clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl b/clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl
index de1b4f8858fa..2bcefe2f192b 100644
--- a/clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl
+++ b/clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl
@@ -6,7 +6,7 @@ global reserve_id_t rid;          // expected-error {{the '__global reserve_id_t
 
 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}}
@@ -16,15 +16,15 @@ 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;
@@ -32,7 +32,7 @@ pipe_int_t test6() {} // expected-error{{declaring function return value of type
 
 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.

diff  --git a/clang/test/SemaOpenCL/null_literal.cl b/clang/test/SemaOpenCL/null_literal.cl
index 2fb287270f46..fb0f938d5952 100644
--- a/clang/test/SemaOpenCL/null_literal.cl
+++ b/clang/test/SemaOpenCL/null_literal.cl
@@ -11,14 +11,14 @@ 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;
 

diff  --git a/clang/test/SemaOpenCL/null_queue.cl b/clang/test/SemaOpenCL/null_queue.cl
index 34d1a621f4a8..a0b33d48dbf7 100644
--- a/clang/test/SemaOpenCL/null_queue.cl
+++ b/clang/test/SemaOpenCL/null_queue.cl
@@ -4,13 +4,13 @@ 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;
 }
@@ -21,7 +21,7 @@ bool compare() {
          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() {

diff  --git a/clang/test/SemaOpenCL/predefined-expr.cl b/clang/test/SemaOpenCL/predefined-expr.cl
index 419e7a925c5c..182c9bda657f 100644
--- a/clang/test/SemaOpenCL/predefined-expr.cl
+++ b/clang/test/SemaOpenCL/predefined-expr.cl
@@ -2,7 +2,7 @@
 // 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__;
 }

diff  --git a/clang/test/SemaOpenCL/queue_t_overload.cl b/clang/test/SemaOpenCL/queue_t_overload.cl
index bc3b241bbe11..9f0cb841180c 100644
--- a/clang/test/SemaOpenCL/queue_t_overload.cl
+++ b/clang/test/SemaOpenCL/queue_t_overload.cl
@@ -1,7 +1,7 @@
 // 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;

diff  --git a/clang/test/SemaOpenCL/shifts.cl b/clang/test/SemaOpenCL/shifts.cl
index 26f59a533ff5..257f70a956d5 100644
--- a/clang/test/SemaOpenCL/shifts.cl
+++ b/clang/test/SemaOpenCL/shifts.cl
@@ -45,7 +45,7 @@ void ptest07() {
 // ** 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) {

diff  --git a/clang/test/SemaOpenCL/to_addr_builtin.cl b/clang/test/SemaOpenCL/to_addr_builtin.cl
index 26389d24fce8..ff2d7807356a 100644
--- a/clang/test/SemaOpenCL/to_addr_builtin.cl
+++ b/clang/test/SemaOpenCL/to_addr_builtin.cl
@@ -13,7 +13,7 @@ void test(void) {
   glob = to_global(glob, loc);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
   // expected-error at -2{{implicit declaration of function 'to_global' is invalid in OpenCL}}
-  // expected-warning at -3{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning at -3{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
   // expected-error at -5{{invalid number of arguments to function: 'to_global'}}
 #endif
@@ -21,46 +21,46 @@ void test(void) {
   int x;
   glob = to_global(x);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
   // expected-error at -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 at -2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
   // expected-error at -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 at -2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
   // expected-error at -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 at -2{{incompatible integer to pointer conversion assigning to '__local int *' from 'int'}}
+  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__local int *__private' from 'int'}}
 #else
-  // expected-error at -4{{assigning '__global int *' to '__local int *' changes address space of pointer}}
+  // expected-error at -4{{assigning '__global int *' to '__local int *__private' changes address space of pointer}}
   // expected-warning at -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 at -2{{implicit declaration of function 'to_private' is invalid in OpenCL}}
-  // expected-warning at -3{{incompatible integer to pointer conversion assigning to '__local int *' from 'int'}}
+  // expected-warning at -3{{incompatible integer to pointer conversion assigning to '__local int *__private' from 'int'}}
 #else
-  // expected-error at -5{{assigning 'int *' to '__local int *' changes address space of pointer}}
+  // expected-error at -5{{assigning '__private int *' to '__local int *__private' changes address space of pointer}}
   // expected-warning at -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 at -2{{implicit declaration of function 'to_local' is invalid in OpenCL}}
-  // expected-warning at -3{{incompatible integer to pointer conversion assigning to '__local int *' from 'int'}}
+  // expected-warning at -3{{incompatible integer to pointer conversion assigning to '__local int *__private' from 'int'}}
   // expected-note at -4{{did you mean 'to_global'}}
   // expected-note at 13{{'to_global' declared here}}
 #else
@@ -69,15 +69,15 @@ void test(void) {
 
   priv = to_global(glob);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning at -2{{incompatible integer to pointer conversion assigning to 'int *' from 'int'}}
+  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__private int *__private' from 'int'}}
 #else
-  // expected-error at -4{{assigning '__global int *' to 'int *' changes address space of pointer}}
+  // expected-error at -4{{assigning '__global int *' to '__private int *__private' changes address space of pointer}}
   // expected-warning at -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 at -2{{incompatible integer to pointer conversion assigning to 'int *' from 'int'}}
+  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__private int *__private' from 'int'}}
 #else
   // expected-warning at -4{{passing non-generic address space pointer to to_private may cause dynamic conversion affecting performance}}
 #endif
@@ -85,48 +85,48 @@ void test(void) {
 
   priv = to_local(glob);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning at -2{{incompatible integer to pointer conversion assigning to 'int *' from 'int'}}
+  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__private int *__private' from 'int'}}
 #else
-  // expected-error at -4{{assigning '__local int *' to 'int *' changes address space of pointer}}
+  // expected-error at -4{{assigning '__local int *' to '__private int *__private' changes address space of pointer}}
   // expected-warning at -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 at -2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
   // expected-warning at -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 at -2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
-  // expected-error at -4{{assigning 'int *' to '__global int *' changes address space of pointer}}
+  // expected-error at -4{{assigning '__private int *' to '__global int *__private' changes address space of pointer}}
   // expected-warning at -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 at -2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__global int *__private' from 'int'}}
 #else
-  // expected-error at -4{{assigning '__local int *' to '__global int *' changes address space of pointer}}
+  // expected-error at -4{{assigning '__local int *' to '__global int *__private' changes address space of pointer}}
   // expected-warning at -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 at -2{{incompatible integer to pointer conversion initializing '__global char *' with an expression of type 'int'}}
+  // expected-warning at -2{{incompatible integer to pointer conversion initializing '__global char *__private' with an expression of type 'int'}}
 #else
-  // expected-warning at -4{{incompatible pointer types initializing '__global char *' with an expression of type '__global int *'}}
+  // expected-warning at -4{{incompatible pointer types initializing '__global char *__private' with an expression of type '__global int *'}}
   // expected-warning at -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 at -2{{incompatible integer to pointer conversion assigning to '__global float *' from 'int'}}
+  // expected-warning at -2{{incompatible integer to pointer conversion assigning to '__global float *__private' from 'int'}}
 #else
-  // expected-warning at -4{{incompatible pointer types assigning to '__global float *' from '__global int *'}}
+  // expected-warning at -4{{incompatible pointer types assigning to '__global float *__private' from '__global int *'}}
   // expected-warning at -5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}
 #endif
 

diff  --git a/clang/test/SemaOpenCL/vec_step.cl b/clang/test/SemaOpenCL/vec_step.cl
index d83ebf11fe03..0a2f6ef30e76 100644
--- a/clang/test/SemaOpenCL/vec_step.cl
+++ b/clang/test/SemaOpenCL/vec_step.cl
@@ -26,7 +26,7 @@ void foo(int3 arg1, int8 arg2) {
   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}}
 }

diff  --git a/clang/test/SemaOpenCL/vector_conv_invalid.cl b/clang/test/SemaOpenCL/vector_conv_invalid.cl
index e32d84fd6a7e..a396809c08c5 100644
--- a/clang/test/SemaOpenCL/vector_conv_invalid.cl
+++ b/clang/test/SemaOpenCL/vector_conv_invalid.cl
@@ -7,7 +7,7 @@ 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)}}
@@ -16,7 +16,7 @@ void vector_conv_invalid(const global int4 *const_global_ptr) {
   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;
 }

diff  --git a/clang/test/SemaOpenCLCXX/address-space-deduction.cl b/clang/test/SemaOpenCLCXX/address-space-deduction.cl
index 9bffeafb1c2d..0275ae57efe4 100644
--- a/clang/test/SemaOpenCLCXX/address-space-deduction.cl
+++ b/clang/test/SemaOpenCLCXX/address-space-deduction.cl
@@ -30,8 +30,8 @@ struct c2 {
 
 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;
@@ -41,8 +41,8 @@ struct x1 {
 
 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;
   }
@@ -57,10 +57,10 @@ void bar(__global x1<int> *xx, __global x2<int> *bar) {
 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) {}
 
@@ -68,7 +68,7 @@ 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 = ⅈ
@@ -92,15 +92,15 @@ __kernel void test() {
 }
 
 // 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() {

diff  --git a/clang/test/SemaOpenCLCXX/address-space-lambda.cl b/clang/test/SemaOpenCLCXX/address-space-lambda.cl
index 7110afaee1ee..eeea71e6353f 100644
--- a/clang/test/SemaOpenCLCXX/address-space-lambda.cl
+++ b/clang/test/SemaOpenCLCXX/address-space-lambda.cl
@@ -1,6 +1,6 @@
 //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() {
@@ -25,16 +25,16 @@ __kernel void test() {
 }
 
 __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}}
@@ -43,7 +43,7 @@ __kernel void test_qual() {
   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}}
 

diff  --git a/clang/test/SemaOpenCLCXX/address-space-templates.cl b/clang/test/SemaOpenCLCXX/address-space-templates.cl
index 0acc5145a632..6b304d2fdda4 100644
--- a/clang/test/SemaOpenCLCXX/address-space-templates.cl
+++ b/clang/test/SemaOpenCLCXX/address-space-templates.cl
@@ -19,7 +19,7 @@ T *foo2(T *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() {

diff  --git a/clang/test/SemaOpenCLCXX/addrspace-auto.cl b/clang/test/SemaOpenCLCXX/addrspace-auto.cl
index 56fd9eb58ddc..2860237ddef7 100644
--- a/clang/test/SemaOpenCLCXX/addrspace-auto.cl
+++ b/clang/test/SemaOpenCLCXX/addrspace-auto.cl
@@ -6,30 +6,30 @@ 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'}}
 }