vendor/clang/clang-trunk-r351319

25 files changed, 817 insertions, 150 deletions

diff --git a/test/SemaOpenCL/access-qualifier.cl b/test/SemaOpenCL/access-qualifier.cl
index 35e838b1bf83..1753a1dbefdd 100644
--- a/test/SemaOpenCL/access-qualifier.cl
+++ b/test/SemaOpenCL/access-qualifier.cl
@@ -60,7 +60,7 @@ kernel void k10(read_only Int img){} // expected-error {{access qualifier can on
 
 kernel void k11(read_only write_only image1d_t i){} // expected-error{{multiple access qualifiers}}
 
-kernel void k12(read_only read_only image1d_t i){} // expected-error{{multiple access qualifiers}}
+kernel void k12(read_only read_only image1d_t i){} // expected-warning {{duplicate 'read_only' declaration specifier}}
 
 #if __OPENCL_C_VERSION__ >= 200
 kernel void k13(read_write pipe int i){} // expected-error{{access qualifier 'read_write' can not be used for 'read_only pipe int'}}
@@ -78,3 +78,33 @@ kernel void k14(read_only pipe int p) {
 #if __OPENCL_C_VERSION__ < 200
 kernel void test_image3d_wo(write_only image3d_t img) {} // expected-error {{use of type '__write_only image3d_t' requires cl_khr_3d_image_writes extension to be enabled}}
 #endif
+
+#if __OPENCL_C_VERSION__ >= 200
+kernel void read_write_twice_typedef(read_write img1d_rw i){} // expected-warning {{duplicate 'read_write' declaration specifier}}
+// expected-note@-74 {{previously declared 'read_write' here}}
+
+kernel void pipe_ro_twice(read_only read_only pipe int i){} // expected-warning{{duplicate 'read_only' declaration specifier}}
+// Conflicting access qualifiers
+kernel void pipe_ro_twice_tw(read_write read_only read_only pipe int i){} // expected-error{{access qualifier 'read_write' can not be used for 'read_only pipe int'}}
+kernel void pipe_ro_wo(read_only write_only pipe int i){} // expected-error{{multiple access qualifiers}}
+
+typedef read_only pipe int ROPipeInt;
+kernel void pipe_ro_twice_typedef(read_only ROPipeInt i){} // expected-warning{{duplicate 'read_only' declaration specifier}}
+// expected-note@-2 {{previously declared 'read_only' here}}
+
+kernel void pass_ro_typedef_to_wo(ROPipeInt p) {
+  myPipeWrite(p); // expected-error {{passing 'ROPipeInt' (aka 'read_only pipe int') to parameter of incompatible type 'write_only pipe int'}}
+  // expected-note@-25 {{passing argument to parameter here}}
+}
+#endif
+
+kernel void read_only_twice_typedef(__read_only img1d_ro i){} // expected-warning {{duplicate '__read_only' declaration specifier}}
+// expected-note@-95 {{previously declared 'read_only' here}}
+
+kernel void read_only_twice_default(read_only img1d_ro_default img){} // expected-warning {{duplicate 'read_only' declaration specifier}}
+// expected-note@-101 {{previously declared 'read_only' here}}
+
+kernel void image_wo_twice(write_only __write_only image1d_t i){} // expected-warning {{duplicate '__write_only' declaration specifier}}
+kernel void image_wo_twice_typedef(write_only img1d_wo i){} // expected-warning {{duplicate 'write_only' declaration specifier}}
+// expected-note@-103 {{previously declared 'write_only' here}}
+
diff --git a/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl b/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
index 73e29e7879d5..619ecc4e4771 100644
--- a/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
+++ b/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
@@ -1,6 +1,9 @@
 // RUN: %clang_cc1 %s -ffake-address-space-map -verify -pedantic -fsyntax-only -DCONSTANT -cl-std=CL2.0
 // RUN: %clang_cc1 %s -ffake-address-space-map -verify -pedantic -fsyntax-only -DGLOBAL -cl-std=CL2.0
 // RUN: %clang_cc1 %s -ffake-address-space-map -verify -pedantic -fsyntax-only -DGENERIC -cl-std=CL2.0
+// RUN: %clang_cc1 %s -ffake-address-space-map -verify -pedantic -fsyntax-only -DCONSTANT -cl-std=c++
+// RUN: %clang_cc1 %s -ffake-address-space-map -verify -pedantic -fsyntax-only -DGLOBAL -cl-std=c++
+// RUN: %clang_cc1 %s -ffake-address-space-map -verify -pedantic -fsyntax-only -DGENERIC -cl-std=c++
 
 /* OpenCLC v2.0 adds a set of restrictions for conversions between pointers to
 *  different address spaces, mainly described in Sections 6.5.5 and 6.5.6.
@@ -17,71 +20,111 @@
 */
 
 #ifdef GENERIC
-#define AS generic
-#define AS_COMP local
-#define AS_INCOMP constant
+#define AS __generic
+#define AS_COMP __local
+#define AS_INCOMP __constant
 #endif
 
 #ifdef GLOBAL
-#define AS global
-#define AS_COMP global
-#define AS_INCOMP local
+#define AS __global
+#define AS_COMP __global
+#define AS_INCOMP __local
 #endif
 
 #ifdef CONSTANT
-#define AS constant
-#define AS_COMP constant
-#define AS_INCOMP global
+#define AS __constant
+#define AS_COMP __constant
+#define AS_INCOMP __global
 #endif
 
-void f_glob(global int *arg_glob) {}
+void f_glob(__global int *arg_glob) {}
 #ifndef GLOBAL
-// expected-note@-2{{passing argument to parameter 'arg_glob' here}}
+#if !__OPENCL_CPP_VERSION__
+// expected-note@-3{{passing argument to parameter 'arg_glob' here}}
+#else
+// expected-note-re@-5{{candidate function not viable: address space mismatch in 1st argument ('__{{generic|constant}} int *'), parameter type must be '__global int *'}}
+#endif
 #endif
 
-void f_loc(local int *arg_loc) {
-} // expected-note@-1{{passing argument to parameter 'arg_loc' here}}
+void f_loc(__local int *arg_loc) {}
+#if !__OPENCL_CPP_VERSION__
+// expected-note@-2{{passing argument to parameter 'arg_loc' here}}
+#else
+// expected-note-re@-4{{candidate function not viable: address space mismatch in 1st argument ('__{{global|generic|constant}} int *'), parameter type must be '__local int *'}}
+#endif
 
-void f_const(constant int *arg_const) {}
+void f_const(__constant int *arg_const) {}
 #ifndef CONSTANT
-// expected-note@-2{{passing argument to parameter 'arg_const' here}}
+#if !__OPENCL_CPP_VERSION__
+// expected-note@-3{{passing argument to parameter 'arg_const' here}}
+#else
+// expected-note-re@-5{{candidate function not viable: address space mismatch in 1st argument ('__{{global|generic}} int *'), parameter type must be '__constant int *'}}
+#endif
 #endif
 
-void f_priv(private int *arg_priv) {
-} // expected-note@-1{{passing argument to parameter 'arg_priv' here}}
+void f_priv(__private int *arg_priv) {}
+#if !__OPENCL_CPP_VERSION__
+// expected-note@-2{{passing argument to parameter 'arg_priv' here}}
+#else
+// expected-note-re@-4{{candidate function not viable: address space mismatch in 1st argument ('__{{global|generic|constant}} int *'), parameter type must be 'int *'}}
+#endif
 
-void f_gen(generic int *arg_gen) {}
+void f_gen(__generic int *arg_gen) {}
 #ifdef CONSTANT
-// expected-note@-2{{passing argument to parameter 'arg_gen' here}}
+#if !__OPENCL_CPP_VERSION__
+// expected-note@-3{{passing argument to parameter 'arg_gen' here}}
+#else
+// expected-note@-5{{candidate function not viable: address space mismatch in 1st argument ('__constant int *'), parameter type must be '__generic int *'}}
+#endif
 #endif
 
-void test_conversion(global int *arg_glob, local int *arg_loc,
-                     constant int *arg_const, private int *arg_priv,
-                     generic int *arg_gen) {
+void test_conversion(__global int *arg_glob, __local int *arg_loc,
+                     __constant int *arg_const, __private int *arg_priv,
+                     __generic int *arg_gen) {
 
   AS int *var_init1 = arg_glob;
 #ifdef CONSTANT
-// expected-error@-2{{initializing '__constant int *' with an expression of type '__global int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{initializing '__constant int *' with an expression of type '__global int *' changes address space of pointer}}
+#else
+// expected-error@-5{{cannot initialize a variable of type '__constant int *' with an lvalue of type '__global int *'}}
+#endif
 #endif
 
   AS int *var_init2 = arg_loc;
 #ifndef GENERIC
-// expected-error-re@-2{{initializing '__{{global|constant}} int *' with an expression of type '__local int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{initializing '__{{global|constant}} int *' with an expression of type '__local int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{cannot initialize a variable of type '__{{global|constant}} int *' with an lvalue of type '__local int *'}}
+#endif
 #endif
 
   AS int *var_init3 = arg_const;
 #ifndef CONSTANT
-// expected-error-re@-2{{initializing '__{{global|generic}} int *' with an expression of type '__constant int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{initializing '__{{global|generic}} int *' with an expression of type '__constant int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{cannot initialize a variable of type '__{{global|generic}} int *' with an lvalue of type '__constant int *'}}
+#endif
 #endif
 
   AS int *var_init4 = arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{initializing '__{{global|constant}} int *' with an expression of type 'int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{initializing '__{{global|constant}} int *' with an expression of type 'int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{cannot initialize a variable of type '__{{global|constant}} int *' with an lvalue of type 'int *'}}
+#endif
 #endif
 
   AS int *var_init5 = arg_gen;
 #ifndef GENERIC
-// expected-error-re@-2{{initializing '__{{global|constant}} int *' with an expression of type '__generic int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{initializing '__{{global|constant}} int *' with an expression of type '__generic int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{cannot initialize a variable of type '__{{global|constant}} int *' with an lvalue of type '__generic int *'}}
+#endif
 #endif
 
   AS int *var_cast1 = (AS int *)arg_glob;
@@ -112,27 +155,47 @@ void test_conversion(global int *arg_glob, local int *arg_loc,
   AS int *var_impl;
   var_impl = arg_glob;
 #ifdef CONSTANT
-// expected-error@-2{{assigning '__global int *' to '__constant int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{assigning '__global int *' to '__constant int *' changes address space of pointer}}
+#else
+// expected-error@-5{{assigning to '__constant int *' from incompatible type '__global int *'}}
+#endif
 #endif
 
   var_impl = arg_loc;
 #ifndef GENERIC
-// expected-error-re@-2{{assigning '__local int *' to '__{{global|constant}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{assigning '__local int *' to '__{{global|constant}} int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type '__local int *'}}
+#endif
 #endif
 
   var_impl = arg_const;
 #ifndef CONSTANT
-// expected-error-re@-2{{assigning '__constant int *' to '__{{global|generic}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{assigning '__constant int *' to '__{{global|generic}} int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{assigning to '__{{global|generic}} int *' from incompatible type '__constant int *'}}
+#endif
 #endif
 
   var_impl = arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{assigning 'int *' to '__{{global|constant}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{assigning 'int *' to '__{{global|constant}} int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type 'int *'}}
+#endif
 #endif
 
   var_impl = arg_gen;
 #ifndef GENERIC
-// expected-error-re@-2{{assigning '__generic int *' to '__{{global|constant}} int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{assigning '__generic int *' to '__{{global|constant}} int *' changes address space of pointer}}
+#else
+// expected-error-re@-5{{assigning to '__{{global|constant}} int *' from incompatible type '__generic int *'}}
+#endif
 #endif
 
   var_cast1 = (AS int *)arg_glob;
@@ -163,27 +226,47 @@ void test_conversion(global int *arg_glob, local int *arg_loc,
   AS int *var_cmp;
   int b = var_cmp != arg_glob;
 #ifdef CONSTANT
-// expected-error@-2{{comparison between  ('__constant int *' and '__global int *') which are pointers to non-overlapping address spaces}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{comparison between  ('__constant int *' and '__global int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error@-5{{comparison of distinct pointer types ('__constant int *' and '__global int *')}}
+#endif
 #endif
 
   b = var_cmp != arg_loc;
 #ifndef GENERIC
-// expected-error-re@-2{{comparison between  ('__{{global|constant}} int *' and '__local int *') which are pointers to non-overlapping address spaces}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{comparison between  ('__{{global|constant}} int *' and '__local int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error-re@-5{{comparison of distinct pointer types ('__{{global|constant}} int *' and '__local int *')}}
+#endif
 #endif
 
   b = var_cmp == arg_const;
 #ifndef CONSTANT
-// expected-error-re@-2{{comparison between  ('__{{global|generic}} int *' and '__constant int *') which are pointers to non-overlapping address spaces}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{comparison between  ('__{{global|generic}} int *' and '__constant int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error-re@-5{{comparison of distinct pointer types ('__{{global|generic}} int *' and '__constant int *')}}
+#endif
 #endif
 
   b = var_cmp <= arg_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{comparison between  ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{comparison between  ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error-re@-5{{comparison of distinct pointer types ('__{{global|constant}} int *' and 'int *')}}
+#endif
 #endif
 
   b = var_cmp >= arg_gen;
 #ifdef CONSTANT
-// expected-error@-2{{comparison between  ('__constant int *' and '__generic int *') which are pointers to non-overlapping address spaces}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{comparison between  ('__constant int *' and '__generic int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error@-5{{comparison of distinct pointer types ('__constant int *' and '__generic int *')}}
+#endif
 #endif
 
   AS int *var_sub;
@@ -214,96 +297,158 @@ void test_conversion(global int *arg_glob, local int *arg_loc,
 
   f_glob(var_sub);
 #ifndef GLOBAL
-// expected-error-re@-2{{passing '__{{constant|generic}} int *' to parameter of type '__global int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{passing '__{{constant|generic}} int *' to parameter of type '__global int *' changes address space of pointer}}
+#else
+// expected-error@-5{{no matching function for call to 'f_glob'}}
+#endif
 #endif
 
-  f_loc(var_sub); // expected-error-re{{passing '__{{global|constant|generic}} int *' to parameter of type '__local int *' changes address space of pointer}}
+  f_loc(var_sub);
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-2{{passing '__{{global|constant|generic}} int *' to parameter of type '__local int *' changes address space of pointer}}
+#else
+// expected-error@-4{{no matching function for call to 'f_loc'}}
+#endif
 
   f_const(var_sub);
 #ifndef CONSTANT
-// expected-error-re@-2{{passing '__{{global|generic}} int *' to parameter of type '__constant int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{passing '__{{global|generic}} int *' to parameter of type '__constant int *' changes address space of pointer}}
+#else
+// expected-error@-5{{no matching function for call to 'f_const'}}
+#endif
 #endif
 
-  f_priv(var_sub); // expected-error-re{{passing '__{{global|constant|generic}} int *' to parameter of type 'int *' changes address space of pointer}}
+  f_priv(var_sub);
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-2{{passing '__{{global|constant|generic}} int *' to parameter of type 'int *' changes address space of pointer}}
+#else
+// expected-error@-4{{no matching function for call to 'f_priv'}}
+#endif
 
   f_gen(var_sub);
 #ifdef CONSTANT
-// expected-error@-2{{passing '__constant int *' to parameter of type '__generic int *' changes address space of pointer}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{passing '__constant int *' to parameter of type '__generic int *' changes address space of pointer}}
+#else
+// expected-error@-5{{no matching function for call to 'f_gen'}}
+#endif
 #endif
 }
 
 void test_ternary() {
   AS int *var_cond;
-  generic int *var_gen;
-  global int *var_glob;
+  __generic int *var_gen;
+  __global int *var_glob;
   var_gen = 0 ? var_cond : var_glob;
 #ifdef CONSTANT
-// expected-error@-2{{conditional operator with the second and third operands of type  ('__constant int *' and '__global int *') which are pointers to non-overlapping address spaces}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{conditional operator with the second and third operands of type  ('__constant int *' and '__global int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error@-5{{incompatible operand types ('__constant int *' and '__global int *')}}
+#endif
 #endif
 
-  local int *var_loc;
+  __local int *var_loc;
   var_gen = 0 ? var_cond : var_loc;
 #ifndef GENERIC
-// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and '__local int *') which are pointers to non-overlapping address spaces}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and '__local int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error-re@-5{{incompatible operand types ('__{{global|constant}} int *' and '__local int *')}}
+#endif
 #endif
 
-  constant int *var_const;
+  __constant int *var_const;
   var_cond = 0 ? var_cond : var_const;
 #ifndef CONSTANT
-// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|generic}} int *' and '__constant int *') which are pointers to non-overlapping address spaces}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{conditional operator with the second and third operands of type  ('__{{global|generic}} int *' and '__constant int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error-re@-5{{incompatible operand types ('__{{global|generic}} int *' and '__constant int *')}}
+#endif
 #endif
 
-  private int *var_priv;
+  __private int *var_priv;
   var_gen = 0 ? var_cond : var_priv;
 #ifndef GENERIC
-// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error-re@-3{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error-re@-5{{incompatible operand types ('__{{global|constant}} int *' and 'int *')}}
+#endif
 #endif
 
   var_gen = 0 ? var_cond : var_gen;
 #ifdef CONSTANT
-// expected-error@-2{{conditional operator with the second and third operands of type  ('__constant int *' and '__generic int *') which are pointers to non-overlapping address spaces}}
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-3{{conditional operator with the second and third operands of type  ('__constant int *' and '__generic int *') which are pointers to non-overlapping address spaces}}
+#else
+// expected-error@-5{{incompatible operand types ('__constant int *' and '__generic int *')}}
+#endif
 #endif
 
   void *var_void_gen;
-  global char *var_glob_ch;
+  __global char *var_glob_ch;
   var_void_gen = 0 ? var_cond : var_glob_ch;
+#if __OPENCL_CPP_VERSION__
+// expected-error-re@-2{{incompatible operand types ('__{{constant|global|generic}} int *' and '__global char *')}}
+#else
 #ifdef CONSTANT
-// expected-error@-2{{conditional operator with the second and third operands of type  ('__constant int *' and '__global char *') which are pointers to non-overlapping address spaces}}
+// expected-error@-5{{conditional operator with the second and third operands of type  ('__constant int *' and '__global char *') which are pointers to non-overlapping address spaces}}
 #else
-// expected-warning-re@-4{{pointer type mismatch ('__{{global|generic}} int *' and '__global char *')}}
+// expected-warning-re@-7{{pointer type mismatch ('__{{global|generic}} int *' and '__global char *')}}
+#endif
 #endif
 
-  local char *var_loc_ch;
+  __local char *var_loc_ch;
   var_void_gen = 0 ? var_cond : var_loc_ch;
+#if __OPENCL_CPP_VERSION__
+// expected-error-re@-2{{incompatible operand types ('__{{constant|global|generic}} int *' and '__local char *')}}
+#else
 #ifndef GENERIC
-// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and '__local char *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-5{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and '__local char *') which are pointers to non-overlapping address spaces}}
 #else
-// expected-warning@-4{{pointer type mismatch ('__generic int *' and '__local char *')}}
+// expected-warning@-7{{pointer type mismatch ('__generic int *' and '__local char *')}}
+#endif
 #endif
 
-  constant void *var_void_const;
-  constant char *var_const_ch;
+  __constant void *var_void_const;
+  __constant char *var_const_ch;
   var_void_const = 0 ? var_cond : var_const_ch;
+#if __OPENCL_CPP_VERSION__
+// expected-error-re@-2{{incompatible operand types ('__{{constant|global|generic}} int *' and '__constant char *')}}
+#else
 #ifndef CONSTANT
-// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|generic}} int *' and '__constant char *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-5{{conditional operator with the second and third operands of type  ('__{{global|generic}} int *' and '__constant char *') which are pointers to non-overlapping address spaces}}
 #else
-// expected-warning@-4{{pointer type mismatch ('__constant int *' and '__constant char *')}}
+// expected-warning@-7{{pointer type mismatch ('__constant int *' and '__constant char *')}}
+#endif
 #endif
 
-  private char *var_priv_ch;
+  __private char *var_priv_ch;
   var_void_gen = 0 ? var_cond : var_priv_ch;
+#if __OPENCL_CPP_VERSION__
+// expected-error-re@-2{{incompatible operand types ('__{{constant|global|generic}} int *' and 'char *')}}
+#else
 #ifndef GENERIC
-// expected-error-re@-2{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and 'char *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-5{{conditional operator with the second and third operands of type  ('__{{global|constant}} int *' and 'char *') which are pointers to non-overlapping address spaces}}
 #else
-// expected-warning@-4{{pointer type mismatch ('__generic int *' and 'char *')}}
+// expected-warning@-7{{pointer type mismatch ('__generic int *' and 'char *')}}
+#endif
 #endif
 
-  generic char *var_gen_ch;
+  __generic char *var_gen_ch;
   var_void_gen = 0 ? var_cond : var_gen_ch;
+#if __OPENCL_CPP_VERSION__
+// expected-error-re@-2{{incompatible operand types ('__{{constant|global|generic}} int *' and '__generic char *')}}
+#else
 #ifdef CONSTANT
-// expected-error@-2{{conditional operator with the second and third operands of type  ('__constant int *' and '__generic char *') which are pointers to non-overlapping address spaces}}
+// expected-error@-5{{conditional operator with the second and third operands of type  ('__constant int *' and '__generic char *') which are pointers to non-overlapping address spaces}}
 #else
-// expected-warning-re@-4{{pointer type mismatch ('__{{global|generic}} int *' and '__generic char *')}}
+// expected-warning-re@-7{{pointer type mismatch ('__{{global|generic}} int *' and '__generic char *')}}
+#endif
 #endif
 }
 
@@ -317,14 +462,26 @@ void test_pointer_chains() {
   //  * address spaces of corresponded most outer pointees overlaps, their canonical types are equal
   //  * CVR, address spaces and canonical types of the rest of pointees are equivalent.
   var_as_as_int = 0 ? var_as_as_int : var_asc_as_int;
-
+#if __OPENCL_CPP_VERSION__
+#ifdef GENERIC
+// expected-error@-3{{incompatible operand types ('__generic int *__generic *' and '__generic int *__local *')}}
+#endif
+#endif
   // Case 2: Corresponded inner pointees has non-overlapping address spaces.
   var_as_as_int = 0 ? var_as_as_int : var_asc_asn_int;
-// expected-warning-re@-1{{pointer type mismatch ('__{{(generic|global|constant)}} int *__{{(generic|global|constant)}} *' and '__{{(local|global|constant)}} int *__{{(constant|local|global)}} *')}}
+#if !__OPENCL_CPP_VERSION__
+// expected-warning-re@-2{{pointer type mismatch ('__{{(generic|global|constant)}} int *__{{(generic|global|constant)}} *' and '__{{(local|global|constant)}} int *__{{(constant|local|global)}} *')}}
+#else
+// expected-error-re@-4{{incompatible operand types ('__{{(generic|global|constant)}} int *__{{(generic|global|constant)}} *' and '__{{(local|global|constant)}} int *__{{(constant|local|global)}} *')}}
+#endif
 
   // Case 3: Corresponded inner pointees has overlapping but not equivalent address spaces.
 #ifdef GENERIC
   var_as_as_int = 0 ? var_as_as_int : var_asc_asc_int;
-// expected-warning-re@-1{{pointer type mismatch ('__{{(generic|global|constant)}} int *__{{(generic|global|constant)}} *' and '__{{(local|global|constant)}} int *__{{(local|global|constant)}} *')}}
+#if !__OPENCL_CPP_VERSION__
+// expected-warning-re@-2{{pointer type mismatch ('__{{(generic|global|constant)}} int *__{{(generic|global|constant)}} *' and '__{{(local|global|constant)}} int *__{{(local|global|constant)}} *')}}
+#else
+// expected-error-re@-4{{incompatible operand types ('__{{generic|global|constant}} int *__{{generic|global|constant}} *' and '__{{local|global|constant}} int *__{{local|global|constant}} *')}}
+#endif
 #endif
 }
diff --git a/test/SemaOpenCL/address-spaces.cl b/test/SemaOpenCL/address-spaces.cl
index 52c891e16b73..30f311d6ef15 100644
--- a/test/SemaOpenCL/address-spaces.cl
+++ b/test/SemaOpenCL/address-spaces.cl
@@ -9,46 +9,47 @@ __kernel void foo(__global int *gip) {
   __local int lj = 2; // expected-error {{'__local' variable cannot have an initializer}}
 
   int *ip;
-// FIXME: Temporarily disable part of the test that doesn't work for C++ yet.
-#if !__OPENCL_CPP_VERSION__
-#if __OPENCL_C_VERSION__ < 200
+#if ((!__OPENCL_CPP_VERSION__) && (__OPENCL_C_VERSION__ < 200))
   ip = gip; // expected-error {{assigning '__global int *' to 'int *' changes address space of pointer}}
   ip = &li; // expected-error {{assigning '__local int *' to 'int *' changes address space of pointer}}
   ip = &ci; // expected-error {{assigning '__constant int *' to 'int *' changes address space of pointer}}
 #else
   ip = gip;
   ip = &li;
-  ip = &ci; // expected-error {{assigning '__constant int *' to '__generic int *' changes address space of pointer}}
+  ip = &ci;
+#if !__OPENCL_CPP_VERSION__
+// expected-error@-2 {{assigning '__constant int *' to '__generic int *' changes address space of pointer}}
+#else
+// expected-error@-4 {{assigning to '__generic int *' from incompatible type '__constant int *'}}
+#endif
 #endif
 }
 
-void explicit_cast(global int* g, local int* l, constant int* c, private int* p, const constant int *cc)
-{
-  g = (global int*) l;    // expected-error {{casting '__local int *' to type '__global int *' changes address space of pointer}}
-  g = (global int*) c;    // expected-error {{casting '__constant int *' to type '__global int *' changes address space of pointer}}
-  g = (global int*) cc;   // expected-error {{casting 'const __constant int *' to type '__global int *' changes address space of pointer}}
-  g = (global int*) p;    // expected-error {{casting 'int *' to type '__global int *' changes address space of pointer}}
+void explicit_cast(__global int *g, __local int *l, __constant int *c, __private int *p, const __constant int *cc) {
+  g = (__global int *)l;  // expected-error {{casting '__local int *' to type '__global int *' changes address space of pointer}}
+  g = (__global int *)c;  // expected-error {{casting '__constant int *' to type '__global int *' changes address space of pointer}}
+  g = (__global int *)cc; // expected-error {{casting 'const __constant int *' to type '__global int *' changes address space of pointer}}
+  g = (__global int *)p;  // expected-error {{casting 'int *' to type '__global int *' changes address space of pointer}}
 
-  l = (local int*) g;     // expected-error {{casting '__global int *' to type '__local int *' changes address space of pointer}}
-  l = (local int*) c;     // expected-error {{casting '__constant int *' to type '__local int *' changes address space of pointer}}
-  l = (local int*) cc;    // expected-error {{casting 'const __constant int *' to type '__local int *' changes address space of pointer}}
-  l = (local int*) p;     // expected-error {{casting 'int *' to type '__local int *' changes address space of pointer}}
+  l = (__local int *)g;  // expected-error {{casting '__global int *' to type '__local int *' changes address space of pointer}}
+  l = (__local int *)c;  // expected-error {{casting '__constant int *' to type '__local int *' changes address space of pointer}}
+  l = (__local int *)cc; // expected-error {{casting 'const __constant int *' to type '__local int *' changes address space of pointer}}
+  l = (__local int *)p;  // expected-error {{casting 'int *' to type '__local int *' changes address space of pointer}}
 
-  c = (constant int*) g;  // expected-error {{casting '__global int *' to type '__constant int *' changes address space of pointer}}
-  c = (constant int*) l;  // expected-error {{casting '__local int *' to type '__constant int *' changes address space of pointer}}
-  c = (constant int*) p;  // expected-error {{casting 'int *' to type '__constant int *' changes address space of pointer}}
+  c = (__constant int *)g; // expected-error {{casting '__global int *' to type '__constant int *' changes address space of pointer}}
+  c = (__constant int *)l; // expected-error {{casting '__local int *' to type '__constant int *' changes address space of pointer}}
+  c = (__constant int *)p; // expected-error {{casting 'int *' to type '__constant int *' changes address space of pointer}}
 
-  p = (private int*) g;   // expected-error {{casting '__global int *' to type 'int *' changes address space of pointer}}
-  p = (private int*) l;   // expected-error {{casting '__local int *' to type 'int *' changes address space of pointer}}
-  p = (private int*) c;   // expected-error {{casting '__constant int *' to type 'int *' changes address space of pointer}}
-  p = (private int*) cc;  // expected-error {{casting 'const __constant int *' to type 'int *' changes address space of pointer}}
+  p = (__private int *)g;  // expected-error {{casting '__global int *' to type 'int *' changes address space of pointer}}
+  p = (__private int *)l;  // expected-error {{casting '__local int *' to type 'int *' changes address space of pointer}}
+  p = (__private int *)c;  // expected-error {{casting '__constant int *' to type 'int *' changes address space of pointer}}
+  p = (__private int *)cc; // expected-error {{casting 'const __constant int *' to type 'int *' changes address space of pointer}}
 }
 
-void ok_explicit_casts(global int *g, global int* g2, local int* l, local int* l2, private int* p, private int* p2)
-{
-  g = (global int*) g2;
-  l = (local int*) l2;
-  p = (private int*) p2;
+void ok_explicit_casts(__global int *g, __global int *g2, __local int *l, __local int *l2, __private int *p, __private int *p2) {
+  g = (__global int *)g2;
+  l = (__local int *)l2;
+  p = (__private int *)p2;
 }
 
 __private int func_return_priv(void);       //expected-error {{return value cannot be qualified with address space}}
@@ -61,11 +62,10 @@ __generic int func_return_generic(void);    //expected-error {{return value cann
 
 void func_multiple_addr(void) {
   typedef __private int private_int_t;
-  __local __private int var1;   // expected-error {{multiple address spaces specified for type}}
-  __local __private int *var2;  // expected-error {{multiple address spaces specified for type}}
+  __private __local int var1;   // expected-error {{multiple address spaces specified for type}}
+  __private __local int *var2;  // expected-error {{multiple address spaces specified for type}}
   __local private_int_t var3;   // expected-error {{multiple address spaces specified for type}}
   __local private_int_t *var4;  // expected-error {{multiple address spaces specified for type}}
   __private private_int_t var5; // expected-warning {{multiple identical address spaces specified for type}}
   __private private_int_t *var6;// expected-warning {{multiple identical address spaces specified for type}}
-#endif // !__OPENCL_CXX_VERSION__
 }
diff --git a/test/SemaOpenCL/atomic-ops.cl b/test/SemaOpenCL/atomic-ops.cl
index 2ec6cde21176..540520d4e325 100644
--- a/test/SemaOpenCL/atomic-ops.cl
+++ b/test/SemaOpenCL/atomic-ops.cl
@@ -58,7 +58,8 @@ void f(atomic_int *i, const atomic_int *ci,
   __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_load(p, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_load(d, memory_order_seq_cst, memory_scope_work_group);
-  __opencl_atomic_load(ci, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+  __opencl_atomic_load(ci, memory_order_seq_cst, memory_scope_work_group);
+  __opencl_atomic_load(i_c, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-constant _Atomic type ('__constant atomic_int *' (aka '__constant _Atomic(int) *') invalid)}}
 
   __opencl_atomic_store(i, 1, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_store(p, 1, memory_order_seq_cst, memory_scope_work_group);
@@ -94,7 +95,7 @@ void f(atomic_int *i, const atomic_int *ci,
 
   __opencl_atomic_init(ci, 0); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
   __opencl_atomic_store(ci, 0, memory_order_release, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
-  __opencl_atomic_load(ci, memory_order_acquire, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+  __opencl_atomic_load(ci, memory_order_acquire, memory_scope_work_group);
 
   __opencl_atomic_init(&gn, 456);
   __opencl_atomic_init(&gn, (void*)0); // expected-warning{{incompatible pointer to integer conversion passing '__generic void *' to parameter of type 'int'}}
diff --git a/test/SemaOpenCL/block-array-capturing.cl b/test/SemaOpenCL/block-array-capturing.cl
new file mode 100644
index 000000000000..e8073fce11e8
--- /dev/null
+++ b/test/SemaOpenCL/block-array-capturing.cl
@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple spir64-unkown-unkown -emit-llvm %s -o -| FileCheck %s
+// expected-no-diagnostics
+
+typedef int (^block_t)();
+
+int block_typedef_kernel(global int* res) {
+  // CHECK: %{{.*}} = alloca <{ i32, i32, i8 addrspace(4)*, [3 x i32] }>
+  int a[3] = {1, 2, 3};
+  // CHECK: call void @llvm.memcpy{{.*}}
+  block_t b = ^() { return a[0]; };
+  return b();
+}
+
+// CHECK: define {{.*}} @__block_typedef_kernel_block_invoke
+// CHECK: %{{.*}} = getelementptr inbounds [3 x i32], [3 x i32] addrspace(4)* %{{.*}}, i64 0, i64 0
+// CHECK-NOT: call void @llvm.memcpy{{.*}}
diff --git a/test/SemaOpenCL/builtins-amdgcn-error-ci.cl b/test/SemaOpenCL/builtins-amdgcn-error-ci.cl
new file mode 100644
index 000000000000..2f656582bee9
--- /dev/null
+++ b/test/SemaOpenCL/builtins-amdgcn-error-ci.cl
@@ -0,0 +1,8 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
+
+void test_ci_biltins()
+{
+  __builtin_amdgcn_s_dcache_inv_vol(); // expected-error {{'__builtin_amdgcn_s_dcache_inv_vol' needs target feature ci-insts}}
+  __builtin_amdgcn_buffer_wbinvl1_vol(); // expected-error {{'__builtin_amdgcn_buffer_wbinvl1_vol' needs target feature ci-insts}}
+}
diff --git a/test/SemaOpenCL/builtins-amdgcn-error-vi.cl b/test/SemaOpenCL/builtins-amdgcn-error-vi.cl
new file mode 100644
index 000000000000..877a04015e56
--- /dev/null
+++ b/test/SemaOpenCL/builtins-amdgcn-error-vi.cl
@@ -0,0 +1,8 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu hawaii -verify -S -o - %s
+
+void test_vi_s_dcache_wb()
+{
+  __builtin_amdgcn_s_dcache_wb(); // expected-error {{'__builtin_amdgcn_s_dcache_wb' needs target feature vi-insts}}
+}
diff --git a/test/SemaOpenCL/builtins-amdgcn-error.cl b/test/SemaOpenCL/builtins-amdgcn-error.cl
index 67c416feacf7..77b9f573adc9 100644
--- a/test/SemaOpenCL/builtins-amdgcn-error.cl
+++ b/test/SemaOpenCL/builtins-amdgcn-error.cl
@@ -102,6 +102,15 @@ void test_mov_dpp2(global int* out, int a, int b, int c, int d, bool e)
   *out = __builtin_amdgcn_mov_dpp(a, 0, 0, 0, e); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
 }
 
+void test_update_dpp2(global int* out, int a, int b, int c, int d, int e, bool f)
+{
+  *out = __builtin_amdgcn_update_dpp(a, b, 0, 0, 0, false);
+  *out = __builtin_amdgcn_update_dpp(a, 0, c, 0, 0, false); // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+  *out = __builtin_amdgcn_update_dpp(a, 0, 0, d, 0, false); // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+  *out = __builtin_amdgcn_update_dpp(a, 0, 0, 0, e, false); // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+  *out = __builtin_amdgcn_update_dpp(a, 0, 0, 0, 0, f); // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
+}
+
 void test_ds_faddf(local float *out, float src, int a) {
   *out = __builtin_amdgcn_ds_faddf(out, src, a, 0, false); // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
   *out = __builtin_amdgcn_ds_faddf(out, src, 0, a, false); // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
diff --git a/test/SemaOpenCL/cl20-device-side-enqueue.cl b/test/SemaOpenCL/cl20-device-side-enqueue.cl
index 207fe7d340ab..8946911c093c 100644
--- a/test/SemaOpenCL/cl20-device-side-enqueue.cl
+++ b/test/SemaOpenCL/cl20-device-side-enqueue.cl
@@ -212,7 +212,7 @@ kernel void work_group_size_tests() {
 
 #pragma OPENCL EXTENSION cl_khr_subgroups : enable
 
-kernel void foo(global int *buf)
+kernel void foo(global unsigned int *buf)
 {
   ndrange_t n;
   buf[0] = get_kernel_max_sub_group_size_for_ndrange(n, ^(){});
@@ -220,7 +220,7 @@ kernel void foo(global int *buf)
   buf[0] = get_kernel_max_sub_group_size_for_ndrange(n, 1); // expected-error{{illegal call to 'get_kernel_max_sub_group_size_for_ndrange', expected block argument type}}
 }
 
-kernel void bar(global int *buf)
+kernel void bar(global unsigned int *buf)
 {
   __private ndrange_t n;
   buf[0] = get_kernel_sub_group_count_for_ndrange(n, ^(){});
@@ -230,13 +230,13 @@ kernel void bar(global int *buf)
 
 #pragma OPENCL EXTENSION cl_khr_subgroups : disable
 
-kernel void foo1(global int *buf)
+kernel void foo1(global unsigned int *buf)
 {
   ndrange_t n;
   buf[0] = get_kernel_max_sub_group_size_for_ndrange(n, ^(){}); // expected-error {{use of declaration 'get_kernel_max_sub_group_size_for_ndrange' requires cl_khr_subgroups extension to be enabled}}
 }
 
-kernel void bar1(global int *buf)
+kernel void bar1(global unsigned int *buf)
 {
   ndrange_t n;
   buf[0] = get_kernel_sub_group_count_for_ndrange(n, ^(){}); // expected-error {{use of declaration 'get_kernel_sub_group_count_for_ndrange' requires cl_khr_subgroups extension to be enabled}}
diff --git a/test/SemaOpenCL/clang-builtin-version.cl b/test/SemaOpenCL/clang-builtin-version.cl
index 270345d86a61..3e270e64c69d 100644
--- a/test/SemaOpenCL/clang-builtin-version.cl
+++ b/test/SemaOpenCL/clang-builtin-version.cl
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -fblocks -verify -pedantic-errors -fsyntax-only -ferror-limit 100
+// RUN: %clang_cc1 %s -fblocks -verify -pedantic -fsyntax-only -ferror-limit 100
 
 // Confirm CL2.0 Clang builtins are not available in earlier versions
 
diff --git a/test/SemaOpenCL/clk_event_t.cl b/test/SemaOpenCL/clk_event_t.cl
new file mode 100644
index 000000000000..b73daf92fa00
--- /dev/null
+++ b/test/SemaOpenCL/clk_event_t.cl
@@ -0,0 +1,25 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0
+
+// Taken from opencl-c.h
+#define CLK_NULL_EVENT (__builtin_astype(((void*)(__SIZE_MAX__)), clk_event_t))
+
+global clk_event_t ce; // expected-error {{the '__global clk_event_t' type cannot be used to declare a program scope variable}}
+
+int clk_event_tests() {
+  event_t e;
+  clk_event_t ce1;
+  clk_event_t ce2;
+
+  if (e == ce1) { // expected-error {{invalid operands to binary expression ('event_t' and 'clk_event_t')}}
+    return 9;
+  }
+
+  if (ce1 != ce2) {
+    return 1;
+  }
+  else if (ce1 == CLK_NULL_EVENT || ce2 != CLK_NULL_EVENT) {
+    return 0;
+  }
+
+  return 2;
+}
diff --git a/test/SemaOpenCL/extension-begin.cl b/test/SemaOpenCL/extension-begin.cl
index 92ea88143233..276e6d7f10b0 100644
--- a/test/SemaOpenCL/extension-begin.cl
+++ b/test/SemaOpenCL/extension-begin.cl
@@ -1,37 +1,29 @@
 // Test this without pch.
-// RUN: %clang_cc1 %s -DHEADER -DHEADER_USER -triple spir-unknown-unknown -verify -pedantic -fsyntax-only
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -verify -pedantic -fsyntax-only
 
 // Test with pch.
-// RUN: %clang_cc1 %s -DHEADER -triple spir-unknown-unknown -emit-pch -o %t -verify -pedantic
-// RUN: %clang_cc1 %s -DHEADER_USER -triple spir-unknown-unknown -include-pch %t -fsyntax-only -verify -pedantic
+// RUN: %clang_cc1 -x cl %S/extension-begin.h -triple spir-unknown-unknown -emit-pch -o %t.pch -pedantic
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -include-pch %t.pch -DIMPLICIT_INCLUDE -DUSE_PCH -fsyntax-only -verify -pedantic
 
-#if defined(HEADER) && !defined(INCLUDED)
-#define INCLUDED 
+// Test with modules
+// RUN: rm -rf %t.modules
+// RUN: mkdir -p %t.modules
+//
+// RUN: %clang_cc1 -cl-std=CL1.2 -DIMPLICIT_INCLUDE -include %S/extension-begin.h -triple spir-unknown-unknown -O0 -emit-llvm -o - -fmodules -fimplicit-module-maps -fmodules-cache-path=%t.modules %s -verify -pedantic
+//
+// RUN: rm -rf %t.modules
+// RUN: mkdir -p %t.modules
+//
+// RUN: %clang_cc1 -cl-std=CL2.0 -DIMPLICIT_INCLUDE -include %S/extension-begin.h -triple spir-unknown-unknown -O0 -emit-llvm -o - -fmodules -fimplicit-module-maps -fmodules-cache-path=%t.modules %s -verify -pedantic
 
-#pragma OPENCL EXTENSION all : begin // expected-warning {{expected 'disable' - ignoring}}
-#pragma OPENCL EXTENSION all : end // expected-warning {{expected 'disable' - ignoring}}
-
-#pragma OPENCL EXTENSION my_ext : begin 
-
-struct A {
-  int a;
-};
-
-typedef struct A TypedefOfA;
-typedef const TypedefOfA* PointerOfA;
-
-void f(void);
-
-__attribute__((overloadable)) void g(long x);
-
-#pragma OPENCL EXTENSION my_ext : end
-#pragma OPENCL EXTENSION my_ext : end // expected-warning {{OpenCL extension end directive mismatches begin directive - ignoring}}
-
-__attribute__((overloadable)) void g(void);
-
-#endif // defined(HEADER) && !defined(INCLUDED)
-
-#ifdef HEADER_USER
+#ifndef IMPLICIT_INCLUDE
+#include "extension-begin.h"
+#endif // IMPLICIT_INCLUDE
+#ifndef USE_PCH
+// expected-warning@extension-begin.h:4 {{expected 'disable' - ignoring}}
+// expected-warning@extension-begin.h:5 {{expected 'disable' - ignoring}}
+// expected-warning@extension-begin.h:21 {{OpenCL extension end directive mismatches begin directive - ignoring}}
+#endif // USE_PCH
 
 #pragma OPENCL EXTENSION my_ext : enable
 void test_f1(void) {
@@ -48,9 +40,7 @@ void test_f2(void) {
   PointerOfA test_A_pointer; // expected-error {{use of type 'PointerOfA' (aka 'const 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@-26 {{candidate disabled due to OpenCL extension}}
-        // expected-note@-22 {{candidate function not viable: requires 0 arguments, but 1 was provided}}
+        // expected-note@extension-begin.h:18 {{candidate unavailable as it requires OpenCL extension 'my_ext' to be enabled}}
+        // expected-note@extension-begin.h:23 {{candidate function not viable: requires 0 arguments, but 1 was provided}}
 }
 
-#endif // HEADER_USER
-
diff --git a/test/SemaOpenCL/extension-begin.h b/test/SemaOpenCL/extension-begin.h
new file mode 100644
index 000000000000..d9865ba0b33a
--- /dev/null
+++ b/test/SemaOpenCL/extension-begin.h
@@ -0,0 +1,26 @@
+#ifndef INCLUDED
+#define INCLUDED
+
+#pragma OPENCL EXTENSION all : begin
+#pragma OPENCL EXTENSION all : end
+
+#pragma OPENCL EXTENSION my_ext : begin
+
+struct A {
+  int a;
+};
+
+typedef struct A TypedefOfA;
+typedef const __private TypedefOfA* PointerOfA;
+
+void f(void);
+
+__attribute__((overloadable)) void g(long x);
+
+#pragma OPENCL EXTENSION my_ext : end
+#pragma OPENCL EXTENSION my_ext : end
+
+__attribute__((overloadable)) void g(void);
+
+#endif // INCLUDED
+
diff --git a/test/SemaOpenCL/extension-version.cl b/test/SemaOpenCL/extension-version.cl
index 714e4c210861..a587f1db99af 100644
--- a/test/SemaOpenCL/extension-version.cl
+++ b/test/SemaOpenCL/extension-version.cl
@@ -283,6 +283,18 @@
 #pragma OPENCL EXTENSION cl_amd_media_ops2: enable
 
 #if (__OPENCL_C_VERSION__ >= 120)
+#ifndef cl_khr_depth_images
+#error "Missing cl_khr_depth_images define"
+#endif
+#else
+#ifdef cl_khr_depth_images
+#error "Incorrect cl_khr_depth_images define"
+#endif
+// expected-warning@+2{{unsupported OpenCL extension 'cl_khr_depth_images' - ignoring}}
+#endif
+#pragma OPENCL EXTENSION cl_khr_depth_images: enable
+
+#if (__OPENCL_C_VERSION__ >= 120)
 #ifndef cl_intel_subgroups
 #error "Missing cl_intel_subgroups define"
 #endif
@@ -300,3 +312,12 @@
 #endif
 #pragma OPENCL EXTENSION cl_intel_subgroups_short : enable
 
+#if (__OPENCL_C_VERSION__ >= 120)
+#ifndef cl_intel_device_side_avc_motion_estimation
+#error "Missing cl_intel_device_side_avc_motion_estimation define"
+#endif
+#else
+// expected-warning@+2{{unsupported OpenCL extension 'cl_intel_device_side_avc_motion_estimation' - ignoring}}
+#endif
+#pragma OPENCL EXTENSION cl_intel_device_side_avc_motion_estimation : enable
+
diff --git a/test/SemaOpenCL/extensions.cl b/test/SemaOpenCL/extensions.cl
index 6afb11e42a6a..5f95e32d4a54 100644
--- a/test/SemaOpenCL/extensions.cl
+++ b/test/SemaOpenCL/extensions.cl
@@ -70,6 +70,13 @@ void f2(void) {
 // expected-error@-2{{use of type 'double' requires cl_khr_fp64 extension to be enabled}}
 #endif
 
+  typedef double double4 __attribute__((ext_vector_type(4)));
+  double4 d4 = {0.0f, 2.0f, 3.0f, 1.0f};
+#ifdef NOFP64
+// expected-error@-3 {{use of type 'double' requires cl_khr_fp64 extension to be enabled}}
+// expected-error@-3 {{use of type 'double4' (vector of 4 'double' values) requires cl_khr_fp64 extension to be enabled}}
+#endif
+
   (void) 1.0;
 
 #ifdef NOFP64
diff --git a/test/SemaOpenCL/format-strings-fixit.cl b/test/SemaOpenCL/format-strings-fixit.cl
new file mode 100644
index 000000000000..b9f949ffe2fc
--- /dev/null
+++ b/test/SemaOpenCL/format-strings-fixit.cl
@@ -0,0 +1,24 @@
+// RUN: cp %s %t
+// RUN: %clang_cc1 -cl-std=CL1.2 -pedantic -Wall -fixit %t
+// RUN: %clang_cc1 -cl-std=CL1.2 -fsyntax-only -pedantic -Wall -Werror %t
+// RUN: %clang_cc1 -cl-std=CL1.2 -E -o - %t | FileCheck %s
+
+typedef __attribute__((ext_vector_type(4))) int int4;
+typedef __attribute__((ext_vector_type(8))) int int8;
+
+int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2)));
+
+
+void vector_fixits() {
+  printf("%v4f", (int4) 123);
+  // CHECK: printf("%v4d", (int4) 123);
+
+  printf("%v8d", (int4) 123);
+  // CHECK: printf("%v4d", (int4) 123);
+
+  printf("%v4d", (int8) 123);
+  // CHECK: printf("%v8d", (int8) 123);
+
+  printf("%v4f", (int8) 123);
+  // CHECK: printf("%v8d", (int8) 123);
+}
diff --git a/test/SemaOpenCL/intel-subgroup-avc-ext-types.cl b/test/SemaOpenCL/intel-subgroup-avc-ext-types.cl
new file mode 100644
index 000000000000..3392e80565f1
--- /dev/null
+++ b/test/SemaOpenCL/intel-subgroup-avc-ext-types.cl
@@ -0,0 +1,105 @@
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=CL1.2 -cl-ext=+cl_intel_device_side_avc_motion_estimation -fsyntax-only -verify %s
+
+#pragma OPENCL EXTENSION cl_intel_device_side_avc_motion_estimation : enable
+
+// All intel_sub_group_avc_* types can only be used as argument or return value
+// of built-in functions defined in the extension.
+// But there are also additional initialization rules:
+//   * All types except intel_sub_group_avc_mce_* types can be initialized with
+//     the corresponding initializer macro defined in opencl-c.h
+//     Currently all these macroses are defined as 0x0
+//   * In previous versions of the extension these macroses was defined as {0},
+//     so initialization with initializer list containing one integer equal to
+//     zero should also work
+
+struct st{};
+// negative test cases for initializers
+void foo(char c, float f, void* v, struct st ss) {
+  intel_sub_group_avc_mce_payload_t payload_mce = 0; // No zero initializer for mce types
+  // expected-error@-1 {{initializing 'intel_sub_group_avc_mce_payload_t' with an expression of incompatible type 'int'}}
+  intel_sub_group_avc_ime_payload_t payload_ime = 1; // No literal initializer for *payload_t types
+  // expected-error@-1 {{initializing 'intel_sub_group_avc_ime_payload_t' with an expression of incompatible type 'int'}}
+  intel_sub_group_avc_ref_payload_t payload_ref = f;
+  // expected-error@-1 {{initializing 'intel_sub_group_avc_ref_payload_t' with an expression of incompatible type 'float'}}
+  intel_sub_group_avc_sic_payload_t payload_sic = ss;
+  // expected-error@-1 {{initializing 'intel_sub_group_avc_sic_payload_t' with an expression of incompatible type 'struct st'}}
+
+  intel_sub_group_avc_mce_result_t result_mce = 0; // No zero initializer for mce types
+  // expected-error@-1 {{initializing 'intel_sub_group_avc_mce_result_t' with an expression of incompatible type 'int'}}
+  intel_sub_group_avc_ime_result_t result_ime = 1; // No literal initializer for *result_t types
+  // expected-error@-1 {{initializing 'intel_sub_group_avc_ime_result_t' with an expression of incompatible type 'int'}}
+  intel_sub_group_avc_ref_result_t result_ref = f;
+  // expected-error@-1 {{initializing 'intel_sub_group_avc_ref_result_t' with an expression of incompatible type 'float'}}
+  intel_sub_group_avc_sic_result_t result_sic = ss;
+  // expected-error@-1 {{initializing 'intel_sub_group_avc_sic_result_t' with an expression of incompatible type 'struct st'}}
+
+  intel_sub_group_avc_ime_result_single_reference_streamout_t sstreamout = v;
+  // expected-error@-1 {{initializing 'intel_sub_group_avc_ime_result_single_reference_streamout_t' with an expression of incompatible type 'void *'}}
+
+  intel_sub_group_avc_ime_result_dual_reference_streamout_t dstreamin_list = {0x0, 0x1};
+  // expected-warning@-1 {{excess elements in struct initializer}}
+  intel_sub_group_avc_ime_dual_reference_streamin_t dstreamin_list2 = {};
+  // expected-error@-1 {{scalar initializer cannot be empty}}
+  intel_sub_group_avc_ime_single_reference_streamin_t dstreamin_list3 = {c};
+  // expected-error@-1 {{initializing 'intel_sub_group_avc_ime_single_reference_streamin_t' with an expression of incompatible type 'char'}}
+  intel_sub_group_avc_ime_dual_reference_streamin_t dstreamin_list4 = {1};
+  // expected-error@-1 {{initializing 'intel_sub_group_avc_ime_dual_reference_streamin_t' with an expression of incompatible type 'int'}}
+}
+
+// negative tests for initializers and assignment
+void far() {
+  intel_sub_group_avc_mce_payload_t payload_mce;
+  intel_sub_group_avc_mce_payload_t payload_mce2 = payload_mce;
+
+  intel_sub_group_avc_ime_payload_t payload_ime;
+  intel_sub_group_avc_ref_payload_t payload_ref = payload_ime;
+  // expected-error@-1 {{initializing 'intel_sub_group_avc_ref_payload_t' with an expression of incompatible type 'intel_sub_group_avc_ime_payload_t'}}
+
+  intel_sub_group_avc_sic_result_t result_sic;
+  intel_sub_group_avc_ime_result_t result_ime;
+  result_sic = result_ime;
+  // expected-error@-1 {{assigning to 'intel_sub_group_avc_sic_result_t' from incompatible type 'intel_sub_group_avc_ime_result_t'}}
+}
+
+// Using 0x0 directly allows us not to include opencl-c.h header and not to
+// redefine all of these CLK_AVC_*_INTITIALIZE_INTEL macro. '0x0' value must
+// be in sync with ones defined in opencl-c.h
+
+// positive test cases
+void bar() {
+  const sampler_t vme_sampler = 0x0;
+
+  intel_sub_group_avc_mce_payload_t payload_mce; // No literal initializer for mce types
+  intel_sub_group_avc_ime_payload_t payload_ime = 0x0;
+  intel_sub_group_avc_ref_payload_t payload_ref = 0x0;
+  intel_sub_group_avc_sic_payload_t payload_sic = 0x0;
+
+  intel_sub_group_avc_mce_result_t result_mce; // No literal initializer for mce types
+  intel_sub_group_avc_ime_result_t result_ime = 0x0;
+  intel_sub_group_avc_ref_result_t result_ref = 0x0;
+  intel_sub_group_avc_sic_result_t result_sic = 0x0;
+
+  intel_sub_group_avc_ime_result_single_reference_streamout_t sstreamout = 0x0;
+  intel_sub_group_avc_ime_result_dual_reference_streamout_t dstreamout = 0x0;
+  intel_sub_group_avc_ime_single_reference_streamin_t sstreamin = 0x0;
+  intel_sub_group_avc_ime_dual_reference_streamin_t dstreamin = 0x0;
+
+  // It is allowed to assign variables of the same types
+  intel_sub_group_avc_mce_payload_t pauload_mce2 = payload_mce;
+
+  // Initialization with initializer list was supported in the first version
+  // of the extension. So we check for backward compatibility here.
+  intel_sub_group_avc_ime_payload_t payload_ime_list = {0};
+  intel_sub_group_avc_ref_payload_t payload_ref_list = {0};
+  intel_sub_group_avc_sic_payload_t payload_sic_list = {0};
+
+  intel_sub_group_avc_ime_result_t result_ime_list = {0};
+  intel_sub_group_avc_ref_result_t result_ref_list = {0};
+  intel_sub_group_avc_sic_result_t result_sic_list = {0};
+
+  intel_sub_group_avc_ime_result_single_reference_streamout_t sstreamout_list = {0};
+  intel_sub_group_avc_ime_result_dual_reference_streamout_t dstreamout_list = {0};
+  intel_sub_group_avc_ime_single_reference_streamin_t sstreamin_list = {0};
+  intel_sub_group_avc_ime_dual_reference_streamin_t dstreamin_list = {0};
+}
+
diff --git a/test/SemaOpenCL/invalid-clk-events-cl2.0.cl b/test/SemaOpenCL/invalid-clk-events-cl2.0.cl
deleted file mode 100644
index 8c8e1c617611..000000000000
--- a/test/SemaOpenCL/invalid-clk-events-cl2.0.cl
+++ /dev/null
@@ -1,3 +0,0 @@
-// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0
-
-global clk_event_t ce; // expected-error {{the '__global clk_event_t' type cannot be used to declare a program scope variable}}
diff --git a/test/SemaOpenCL/invalid-kernel-attrs.cl b/test/SemaOpenCL/invalid-kernel-attrs.cl
index a19a989f41bb..a96e85c9ee24 100644
--- a/test/SemaOpenCL/invalid-kernel-attrs.cl
+++ b/test/SemaOpenCL/invalid-kernel-attrs.cl
@@ -37,3 +37,10 @@ kernel __attribute__((reqd_work_group_size(0,1,2))) void kernel13(){} // expecte
 __attribute__((intel_reqd_sub_group_size(8))) void kernel14(){} // expected-error {{attribute 'intel_reqd_sub_group_size' can only be applied to an OpenCL kernel}}
 kernel __attribute__((intel_reqd_sub_group_size(0))) void kernel15(){} // expected-error {{'intel_reqd_sub_group_size' attribute must be greater than 0}}
 kernel __attribute__((intel_reqd_sub_group_size(8))) __attribute__((intel_reqd_sub_group_size(16))) void kernel16() {}  //expected-warning{{attribute 'intel_reqd_sub_group_size' is already applied with different parameters}}
+
+__kernel __attribute__((work_group_size_hint(8,-16,32))) void neg1() {} //expected-error{{'work_group_size_hint' attribute requires a non-negative integral compile time constant expression}}
+__kernel __attribute__((reqd_work_group_size(8,16,-32))) void neg2(){} // expected-error{{'reqd_work_group_size' attribute requires a non-negative integral compile time constant expression}}
+
+// 4294967294 is a negative integer if treated as signed.
+// Should compile successfully, since we expect an unsigned.
+__kernel __attribute__((reqd_work_group_size(8,16,4294967294))) void ok1(){}
diff --git a/test/SemaOpenCL/null_queue.cl b/test/SemaOpenCL/null_queue.cl
index 518f7138de6c..34d1a621f4a8 100644
--- a/test/SemaOpenCL/null_queue.cl
+++ b/test/SemaOpenCL/null_queue.cl
@@ -1,12 +1,30 @@
 // RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only
 extern queue_t get_default_queue();
 
-bool compare() {
-  return 1 == get_default_queue() && // expected-error{{invalid operands to binary expression ('int' and 'queue_t')}}
-         get_default_queue() == 1; // expected-error{{invalid operands to binary expression ('queue_t' and 'int')}}
-}
+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 q = 0;
 }
+
+void assign() {
+  queue_t q2, q3;
+  q2 = 5; // expected-error{{assigning to 'queue_t' from incompatible type 'int'}}
+  q3 = 0;
+  q2 = q3 = 0;
+}
+
+bool compare() {
+  queue_t q4, q5;
+  return 1 == get_default_queue() && // expected-error{{invalid operands to binary expression ('int' and 'queue_t')}}
+         get_default_queue() == 1 && // expected-error{{invalid operands to binary expression ('queue_t' and 'int')}}
+	     q4 == q5 &&
+	     q4 != 0 &&
+	     q4 != 0.0f; // expected-error{{invalid operands to binary expression ('queue_t' and 'float')}}
+}
+
+void call() {
+  queue_arg(5); // expected-error {{passing 'int' to parameter of incompatible type 'queue_t'}}
+  queue_arg(0);
+}
diff --git a/test/SemaOpenCL/numbered-address-space.cl b/test/SemaOpenCL/numbered-address-space.cl
new file mode 100644
index 000000000000..423d03274cef
--- /dev/null
+++ b/test/SemaOpenCL/numbered-address-space.cl
@@ -0,0 +1,31 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -verify -pedantic -fsyntax-only %s
+
+void test_numeric_as_to_generic_implicit_cast(__attribute__((address_space(3))) int *as3_ptr, float src) {
+  generic int* generic_ptr = as3_ptr; // FIXME: This should error
+}
+
+void test_numeric_as_to_generic_explicit_cast(__attribute__((address_space(3))) int *as3_ptr, float src) {
+  generic int* generic_ptr = (generic int*) as3_ptr; // Should maybe be valid?
+}
+
+void test_generic_to_numeric_as_implicit_cast() {
+  generic int* generic_ptr = 0;
+  __attribute__((address_space(3))) int *as3_ptr = generic_ptr; // expected-error{{initializing '__attribute__((address_space(3))) int *' with an expression of type '__generic int *' changes address space of pointer}}
+}
+
+void test_generic_to_numeric_as_explicit_cast() {
+  generic int* generic_ptr = 0;
+  __attribute__((address_space(3))) int *as3_ptr = (__attribute__((address_space(3))) int *)generic_ptr;
+}
+
+void test_generic_as_to_builtin_parameter_explicit_cast_numeric(__attribute__((address_space(3))) int *as3_ptr, float src) {
+  generic int* generic_ptr = as3_ptr; // FIXME: This should error
+  volatile float result = __builtin_amdgcn_ds_fmaxf((__attribute__((address_space(3))) float*) generic_ptr, src, 0, 0, false); // expected-error {{passing '__attribute__((address_space(3))) float *' to parameter of type '__local float *' changes address space of pointer}}
+}
+
+void test_generic_as_to_builtin_parameterimplicit_cast_numeric(__attribute__((address_space(3))) int *as3_ptr, float src) {
+  generic int* generic_ptr = as3_ptr;
+  volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false); // expected-warning {{incompatible pointer types passing '__generic int *' to parameter of type '__local float *'}}
+}
+
diff --git a/test/SemaOpenCL/printf-format-string-warnings.cl b/test/SemaOpenCL/printf-format-string-warnings.cl
new file mode 100644
index 000000000000..2b9c5cc3f319
--- /dev/null
+++ b/test/SemaOpenCL/printf-format-string-warnings.cl
@@ -0,0 +1,13 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0 -finclude-default-header
+
+// Make sure warnings are produced based on printf format strings.
+
+
+kernel void format_string_warnings(__constant char* arg) {
+
+  printf("%d", arg); // expected-warning {{format specifies type 'int' but the argument has type '__constant char *'}}
+
+  printf("not enough arguments %d %d", 4); // expected-warning {{more '%' conversions than data arguments}}
+
+  printf("too many arguments", 4); // expected-warning {{data argument not used by format string}}
+}
diff --git a/test/SemaOpenCL/printf-format-strings.cl b/test/SemaOpenCL/printf-format-strings.cl
new file mode 100644
index 000000000000..079a83495685
--- /dev/null
+++ b/test/SemaOpenCL/printf-format-strings.cl
@@ -0,0 +1,94 @@
+// RUN: %clang_cc1 -cl-std=CL1.2 -cl-ext=+cl_khr_fp64 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -cl-std=CL1.2 -cl-ext=-cl_khr_fp64 -fsyntax-only -verify %s
+
+typedef __attribute__((ext_vector_type(2))) float float2;
+typedef __attribute__((ext_vector_type(4))) float float4;
+
+typedef __attribute__((ext_vector_type(2))) int int2;
+typedef __attribute__((ext_vector_type(4))) int int4;
+typedef __attribute__((ext_vector_type(16))) int int16;
+
+int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2)));
+
+kernel void format_v4f32(float4 arg)
+{
+#ifdef cl_khr_fp64
+    printf("%v4f\n", arg);
+
+    // Precision modifier
+    printf("%.2v4f\n", arg);
+#else
+    // FIXME: These should not warn, and the type should be expected to be float.
+    printf("%v4f\n", arg);  // expected-warning {{double __attribute__((ext_vector_type(4)))' but the argument has type 'float4' (vector of 4 'float' values)}}
+
+    // Precision modifier
+    printf("%.2v4f\n", arg); // expected-warning {{double __attribute__((ext_vector_type(4)))' but the argument has type 'float4' (vector of 4 'float' values)}}
+#endif
+}
+
+kernel void format_only_v(int arg)
+{
+    printf("%v", arg); // expected-warning {{incomplete format specifier}}
+}
+
+kernel void format_missing_num(int arg)
+{
+    printf("%v4", arg); // expected-warning {{incomplete format specifier}}
+}
+
+kernel void format_not_num(int arg)
+{
+    printf("%vNd", arg); // expected-warning {{incomplete format specifier}}
+    printf("%v*d", arg); // expected-warning {{incomplete format specifier}}
+}
+
+kernel void format_v16i32(int16 arg)
+{
+    printf("%v16d\n", arg);
+}
+
+kernel void format_v4i32_scalar(int arg)
+{
+   printf("%v4d\n", arg); // expected-warning  {{format specifies type 'int __attribute__((ext_vector_type(4)))' but the argument has type 'int'}}
+}
+
+kernel void format_v4i32_wrong_num_elts_2_to_4(int2 arg)
+{
+    printf("%v4d\n", arg); // expected-warning {{format specifies type 'int __attribute__((ext_vector_type(4)))' but the argument has type 'int2' (vector of 2 'int' values)}}
+}
+
+kernel void format_missing_num_elts_format(int4 arg)
+{
+    printf("%vd\n", arg); // expected-warning {{incomplete format specifier}}
+}
+
+kernel void format_v4f32_scalar(float arg)
+{
+    printf("%v4f\n", arg); // expected-warning {{format specifies type 'double __attribute__((ext_vector_type(4)))' but the argument has type 'float'}}
+}
+
+kernel void format_v4f32_wrong_num_elts(float2 arg)
+{
+    printf("%v4f\n", arg); // expected-warning {{format specifies type 'double __attribute__((ext_vector_type(4)))' but the argument has type 'float2' (vector of 2 'float' values)}}
+}
+
+kernel void format_missing_num_elts(float4 arg)
+{
+    printf("%vf\n", arg); // expected-warning {{incomplete format specifier}}
+}
+
+kernel void vector_precision_modifier_v4i32_to_v4f32(int4 arg)
+{
+    printf("%.2v4f\n", arg); // expected-warning {{format specifies type 'double __attribute__((ext_vector_type(4)))' but the argument has type 'int4' (vector of 4 'int' values)}}
+}
+
+kernel void invalid_Y(int4 arg)
+{
+    printf("%v4Y\n", arg); // expected-warning {{invalid conversion specifier 'Y'}}
+}
+
+// FIXME: This should warn
+kernel void crash_on_s(int4 arg)
+{
+    printf("%v4s\n", arg);
+}
diff --git a/test/SemaOpenCL/sampler_t.cl b/test/SemaOpenCL/sampler_t.cl
index 5a1850b02cbd..8473fa33631a 100644
--- a/test/SemaOpenCL/sampler_t.cl
+++ b/test/SemaOpenCL/sampler_t.cl
@@ -33,6 +33,8 @@ constant sampler_t glb_smp9 = 0x100000000LL; // expected-error{{sampler_t initia
 
 void foo(sampler_t); // expected-note{{passing argument to parameter here}}
 
+void constant_sampler(constant sampler_t s); // expected-error{{parameter may not be qualified with an address space}}
+
 constant struct sampler_s {
   sampler_t smp; // expected-error{{the 'sampler_t' type cannot be used to declare a structure or union field}}
 } sampler_str = {0};
diff --git a/test/SemaOpenCL/to_addr_builtin.cl b/test/SemaOpenCL/to_addr_builtin.cl
index 70956f007a5f..26389d24fce8 100644
--- a/test/SemaOpenCL/to_addr_builtin.cl
+++ b/test/SemaOpenCL/to_addr_builtin.cl
@@ -1,16 +1,18 @@
 // RUN: %clang_cc1 -verify -fsyntax-only %s
-// RUN: %clang_cc1 -verify -fsyntax-only -cl-std=CL2.0 %s
+// RUN: %clang_cc1 -Wconversion -verify -fsyntax-only -cl-std=CL2.0 %s
 
 void test(void) {
   global int *glob;
   local int *loc;
   constant int *con;
+  private int *priv;
+  global float *glob_wrong_ty;
   typedef constant int const_int_ty;
   const_int_ty *con_typedef;
 
   glob = to_global(glob, loc);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-warning@-2{{implicit declaration of function 'to_global' is invalid in OpenCL}}
+  // expected-error@-2{{implicit declaration of function 'to_global' is invalid in OpenCL}}
   // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
 #else
   // expected-error@-5{{invalid number of arguments to function: 'to_global'}}
@@ -43,6 +45,73 @@ void test(void) {
   // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__local int *' from 'int'}}
 #else
   // expected-error@-4{{assigning '__global int *' to '__local int *' changes address space of pointer}}
+  // expected-warning@-5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}
+#endif
+
+  loc = to_private(glob);
+#if __OPENCL_C_VERSION__ < CL_VERSION_2_0
+  // expected-error@-2{{implicit declaration of function 'to_private' is invalid in OpenCL}}
+  // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__local int *' from 'int'}}
+#else
+  // expected-error@-5{{assigning 'int *' to '__local int *' changes address space of pointer}}
+  // expected-warning@-6{{passing non-generic address space pointer to to_private may cause dynamic conversion affecting performance}}
+#endif
+
+  loc = to_local(glob);
+#if __OPENCL_C_VERSION__ < CL_VERSION_2_0
+  // expected-error@-2{{implicit declaration of function 'to_local' is invalid in OpenCL}}
+  // expected-warning@-3{{incompatible integer to pointer conversion assigning to '__local int *' from 'int'}}
+  // expected-note@-4{{did you mean 'to_global'}}
+  // expected-note@13{{'to_global' declared here}}
+#else
+  // expected-warning@-7{{passing non-generic address space pointer to to_local may cause dynamic conversion affecting performance}}
+#endif
+
+  priv = to_global(glob);
+#if __OPENCL_C_VERSION__ < CL_VERSION_2_0
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to 'int *' from 'int'}}
+#else
+  // expected-error@-4{{assigning '__global int *' to 'int *' changes address space of pointer}}
+  // expected-warning@-5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}
+#endif
+
+  priv = to_private(glob);
+#if __OPENCL_C_VERSION__ < CL_VERSION_2_0
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to 'int *' from 'int'}}
+#else
+  // expected-warning@-4{{passing non-generic address space pointer to to_private may cause dynamic conversion affecting performance}}
+#endif
+
+
+  priv = to_local(glob);
+#if __OPENCL_C_VERSION__ < CL_VERSION_2_0
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to 'int *' from 'int'}}
+#else
+  // expected-error@-4{{assigning '__local int *' to 'int *' changes address space of pointer}}
+  // expected-warning@-5{{passing non-generic address space pointer to to_local may cause dynamic conversion affecting performance}}
+#endif
+
+  glob = to_global(glob);
+#if __OPENCL_C_VERSION__ < CL_VERSION_2_0
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+#else
+  // expected-warning@-4{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}
+#endif
+
+  glob = to_private(glob);
+#if __OPENCL_C_VERSION__ < CL_VERSION_2_0
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+#else
+  // expected-error@-4{{assigning 'int *' to '__global int *' changes address space of pointer}}
+  // expected-warning@-5{{passing non-generic address space pointer to to_private may cause dynamic conversion affecting performance}}
+#endif
+
+  glob = to_local(glob);
+#if __OPENCL_C_VERSION__ < CL_VERSION_2_0
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global int *' from 'int'}}
+#else
+  // expected-error@-4{{assigning '__local int *' to '__global int *' changes address space of pointer}}
+  // expected-warning@-5{{passing non-generic address space pointer to to_local may cause dynamic conversion affecting performance}}
 #endif
 
   global char *glob_c = to_global(loc);
@@ -50,6 +119,15 @@ void test(void) {
   // expected-warning@-2{{incompatible integer to pointer conversion initializing '__global char *' with an expression of type 'int'}}
 #else
   // expected-warning@-4{{incompatible pointer types initializing '__global char *' with an expression of type '__global int *'}}
+  // expected-warning@-5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}
+#endif
+
+  glob_wrong_ty = to_global(glob);
+#if __OPENCL_C_VERSION__ < CL_VERSION_2_0
+  // expected-warning@-2{{incompatible integer to pointer conversion assigning to '__global float *' from 'int'}}
+#else
+  // expected-warning@-4{{incompatible pointer types assigning to '__global float *' from '__global int *'}}
+  // expected-warning@-5{{passing non-generic address space pointer to to_global may cause dynamic conversion affecting performance}}
 #endif
 
 }