vendor/clang/clang-trunk-r321017

16 files changed, 353 insertions, 54 deletions

diff --git a/test/SemaOpenCL/address-spaces.cl b/test/SemaOpenCL/address-spaces.cl
index 7026d1bc086d6..97dd1f4f90b3e 100644
--- a/test/SemaOpenCL/address-spaces.cl
+++ b/test/SemaOpenCL/address-spaces.cl
@@ -1,4 +1,5 @@
 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only
 
 __constant int ci = 1;
 
@@ -7,9 +8,15 @@ __kernel void foo(__global int *gip) {
   __local int lj = 2; // expected-error {{'__local' variable cannot have an initializer}}
 
   int *ip;
+#if __OPENCL_C_VERSION__ < 200
   ip = gip; // expected-error {{assigning '__global int *' to 'int *' changes address space of pointer}}
   ip = &li; // expected-error {{assigning '__local int *' to 'int *' changes address space of pointer}}
   ip = &ci; // expected-error {{assigning '__constant int *' to 'int *' changes address space of pointer}}
+#else
+  ip = gip;
+  ip = &li;
+  ip = &ci; // expected-error {{assigning '__constant int *' to '__generic int *' changes address space of pointer}}
+#endif
 }
 
 void explicit_cast(global int* g, local int* l, constant int* c, private int* p, const constant int *cc)
@@ -40,3 +47,19 @@ void ok_explicit_casts(global int *g, global int* g2, local int* l, local int* l
   l = (local int*) l2;
   p = (private int*) p2;
 }
+
+__private int func_return_priv(void);       //expected-error {{return value cannot be qualified with address space}}
+__global int func_return_global(void);      //expected-error {{return value cannot be qualified with address space}}
+__local int func_return_local(void);        //expected-error {{return value cannot be qualified with address space}}
+__constant int func_return_constant(void);  //expected-error {{return value cannot be qualified with address space}}
+#if __OPENCL_C_VERSION__ >= 200
+__generic int func_return_generic(void);    //expected-error {{return value cannot be qualified with address space}}
+#endif
+
+void func_multiple_addr(void) {
+  typedef __private int private_int_t;
+  __local __private int var1;   // expected-error {{multiple address spaces specified for type}}
+  __local __private 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}}
+}
diff --git a/test/SemaOpenCL/atomic-ops.cl b/test/SemaOpenCL/atomic-ops.cl
new file mode 100644
index 0000000000000..2ec6cde211760
--- /dev/null
+++ b/test/SemaOpenCL/atomic-ops.cl
@@ -0,0 +1,195 @@
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -fsyntax-only -triple=spir64
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -fsyntax-only -triple=amdgcn-amdhsa-amd-opencl
+
+// Basic parsing/Sema tests for __opencl_atomic_*
+
+#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
+#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable
+
+typedef __INTPTR_TYPE__ intptr_t;
+typedef int int8 __attribute__((ext_vector_type(8)));
+
+typedef enum memory_order {
+  memory_order_relaxed = __ATOMIC_RELAXED,
+  memory_order_acquire = __ATOMIC_ACQUIRE,
+  memory_order_release = __ATOMIC_RELEASE,
+  memory_order_acq_rel = __ATOMIC_ACQ_REL,
+  memory_order_seq_cst = __ATOMIC_SEQ_CST
+} memory_order;
+
+typedef enum memory_scope {
+  memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
+  memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
+  memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
+  memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
+#if defined(cl_intel_subgroups) || defined(cl_khr_subgroups)
+  memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
+#endif
+} memory_scope;
+
+struct S { char c[3]; };
+
+char i8;
+short i16;
+int i32;
+int8 i64;
+
+atomic_int gn;
+void f(atomic_int *i, const atomic_int *ci,
+       atomic_intptr_t *p, atomic_float *d,
+       int *I, const int *CI,
+       intptr_t *P, float *D, struct S *s1, struct S *s2,
+       global atomic_int *i_g, local atomic_int *i_l, private atomic_int *i_p,
+       constant atomic_int *i_c) {
+  __opencl_atomic_init(I, 5); // expected-error {{address argument to atomic operation must be a pointer to _Atomic type ('__generic int *' invalid)}}
+  __opencl_atomic_init(ci, 5); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+
+  __opencl_atomic_load(0); // expected-error {{too few arguments to function call, expected 3, have 1}}
+  __opencl_atomic_load(0, 0, 0, 0); // expected-error {{too many arguments to function call, expected 3, have 4}}
+  __opencl_atomic_store(0,0,0,0); // expected-error {{address argument to atomic builtin must be a pointer}}
+  __opencl_atomic_store((int *)0, 0, 0, 0); // expected-error {{address argument to atomic operation must be a pointer to _Atomic type ('__generic int *' invalid)}}
+  __opencl_atomic_store(i, 0, memory_order_relaxed, memory_scope_work_group);
+  __opencl_atomic_store(ci, 0, memory_order_relaxed, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+  __opencl_atomic_store(i_g, 0, memory_order_relaxed, memory_scope_work_group);
+  __opencl_atomic_store(i_l, 0, memory_order_relaxed, memory_scope_work_group);
+  __opencl_atomic_store(i_p, 0, memory_order_relaxed, memory_scope_work_group);
+  __opencl_atomic_store(i_c, 0, memory_order_relaxed, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-constant _Atomic type ('__constant atomic_int *' (aka '__constant _Atomic(int) *') invalid)}}
+
+  __opencl_atomic_load(i, memory_order_seq_cst, memory_scope_work_group);
+  __opencl_atomic_load(p, memory_order_seq_cst, memory_scope_work_group);
+  __opencl_atomic_load(d, memory_order_seq_cst, memory_scope_work_group);
+  __opencl_atomic_load(ci, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+
+  __opencl_atomic_store(i, 1, memory_order_seq_cst, memory_scope_work_group);
+  __opencl_atomic_store(p, 1, memory_order_seq_cst, memory_scope_work_group);
+  (int)__opencl_atomic_store(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{operand of type 'void' where arithmetic or pointer type is required}}
+
+  int exchange_1 = __opencl_atomic_exchange(i, 1, memory_order_seq_cst, memory_scope_work_group);
+  int exchange_2 = __opencl_atomic_exchange(I, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to _Atomic}}
+
+  __opencl_atomic_fetch_add(i, 1, memory_order_seq_cst, memory_scope_work_group);
+  __opencl_atomic_fetch_add(p, 1, memory_order_seq_cst, memory_scope_work_group);
+  __opencl_atomic_fetch_add(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
+  __opencl_atomic_fetch_and(i, 1, memory_order_seq_cst, memory_scope_work_group);
+  __opencl_atomic_fetch_and(p, 1, memory_order_seq_cst, memory_scope_work_group);
+  __opencl_atomic_fetch_and(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to bitwise atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
+
+  __opencl_atomic_fetch_min(i, 1, memory_order_seq_cst, memory_scope_work_group);
+  __opencl_atomic_fetch_max(i, 1, memory_order_seq_cst, memory_scope_work_group);
+  __opencl_atomic_fetch_min(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
+  __opencl_atomic_fetch_max(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
+
+  bool cmpexch_1 = __opencl_atomic_compare_exchange_strong(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
+  bool cmpexch_2 = __opencl_atomic_compare_exchange_strong(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
+  bool cmpexch_3 = __opencl_atomic_compare_exchange_strong(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *' to parameter of type '__generic float *'}}
+  (void)__opencl_atomic_compare_exchange_strong(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *' to parameter of type '__generic int *' discards qualifiers}}
+
+  bool 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}}
+
+  // Pointers to different address spaces are allowed.
+  bool cmpexch_10 = __opencl_atomic_compare_exchange_strong((global atomic_int *)0x308, (constant int *)0x309, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
+
+  __opencl_atomic_init(ci, 0); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+  __opencl_atomic_store(ci, 0, memory_order_release, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+  __opencl_atomic_load(ci, memory_order_acquire, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+
+  __opencl_atomic_init(&gn, 456);
+  __opencl_atomic_init(&gn, (void*)0); // expected-warning{{incompatible pointer to integer conversion passing '__generic void *' to parameter of type 'int'}}
+}
+
+void memory_checks(atomic_int *Ap, int *p, int val) {
+  // non-integer memory order argument is casted to integer type.
+  (void)__opencl_atomic_load(Ap, 1.0f, memory_scope_work_group);
+  float forder;
+  (void)__opencl_atomic_load(Ap, forder, memory_scope_work_group);
+  struct S s;
+  (void)__opencl_atomic_load(Ap, s, memory_scope_work_group); // expected-error {{passing 'struct S' to parameter of incompatible type 'int'}}
+
+  (void)__opencl_atomic_load(Ap, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_load(Ap, memory_order_acquire, memory_scope_work_group);
+  (void)__opencl_atomic_load(Ap, memory_order_consume, memory_scope_work_group); // expected-error {{use of undeclared identifier 'memory_order_consume'}}
+  (void)__opencl_atomic_load(Ap, memory_order_release, memory_scope_work_group); // expected-warning {{memory order argument to atomic operation is invalid}}
+  (void)__opencl_atomic_load(Ap, memory_order_acq_rel, memory_scope_work_group); // expected-warning {{memory order argument to atomic operation is invalid}}
+  (void)__opencl_atomic_load(Ap, memory_order_seq_cst, memory_scope_work_group);
+
+  (void)__opencl_atomic_store(Ap, val, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_store(Ap, val, memory_order_acquire, memory_scope_work_group); // expected-warning {{memory order argument to atomic operation is invalid}}
+  (void)__opencl_atomic_store(Ap, val, memory_order_release, memory_scope_work_group);
+  (void)__opencl_atomic_store(Ap, val, memory_order_acq_rel, memory_scope_work_group); // expected-warning {{memory order argument to atomic operation is invalid}}
+  (void)__opencl_atomic_store(Ap, val, memory_order_seq_cst, memory_scope_work_group);
+
+  (void)__opencl_atomic_fetch_add(Ap, 1, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_add(Ap, 1, memory_order_acquire, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_add(Ap, 1, memory_order_release, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_add(Ap, 1, memory_order_acq_rel, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_add(Ap, 1, memory_order_seq_cst, memory_scope_work_group);
+
+  (void)__opencl_atomic_init(Ap, val);
+
+  (void)__opencl_atomic_fetch_sub(Ap, val, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_sub(Ap, val, memory_order_acquire, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_sub(Ap, val, memory_order_release, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_sub(Ap, val, memory_order_acq_rel, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_sub(Ap, val, memory_order_seq_cst, memory_scope_work_group);
+
+  (void)__opencl_atomic_fetch_and(Ap, val, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_and(Ap, val, memory_order_acquire, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_and(Ap, val, memory_order_release, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_and(Ap, val, memory_order_acq_rel, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_and(Ap, val, memory_order_seq_cst, memory_scope_work_group);
+
+  (void)__opencl_atomic_fetch_or(Ap, val, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_or(Ap, val, memory_order_acquire, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_or(Ap, val, memory_order_release, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_or(Ap, val, memory_order_acq_rel, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_or(Ap, val, memory_order_seq_cst, memory_scope_work_group);
+
+  (void)__opencl_atomic_fetch_xor(Ap, val, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_xor(Ap, val, memory_order_acquire, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_xor(Ap, val, memory_order_release, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_xor(Ap, val, memory_order_acq_rel, memory_scope_work_group);
+  (void)__opencl_atomic_fetch_xor(Ap, val, memory_order_seq_cst, memory_scope_work_group);
+
+  (void)__opencl_atomic_exchange(Ap, val, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_exchange(Ap, val, memory_order_acquire, memory_scope_work_group);
+  (void)__opencl_atomic_exchange(Ap, val, memory_order_release, memory_scope_work_group);
+  (void)__opencl_atomic_exchange(Ap, val, memory_order_acq_rel, memory_scope_work_group);
+  (void)__opencl_atomic_exchange(Ap, val, memory_order_seq_cst, memory_scope_work_group);
+
+  (void)__opencl_atomic_compare_exchange_strong(Ap, p, val, memory_order_relaxed, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_compare_exchange_strong(Ap, p, val, memory_order_acquire, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_compare_exchange_strong(Ap, p, val, memory_order_release, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_compare_exchange_strong(Ap, p, val, memory_order_acq_rel, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, memory_order_relaxed, memory_scope_work_group);
+
+  (void)__opencl_atomic_compare_exchange_weak(Ap, p, val, memory_order_relaxed, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_compare_exchange_weak(Ap, p, val, memory_order_acquire, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_compare_exchange_weak(Ap, p, val, memory_order_release, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_compare_exchange_weak(Ap, p, val, memory_order_acq_rel, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, memory_order_relaxed, memory_scope_work_group);
+}
+
+void synchscope_checks(atomic_int *Ap, int scope) {
+  (void)__opencl_atomic_load(Ap, memory_order_relaxed, memory_scope_work_item); // expected-error{{synchronization scope argument to atomic operation is invalid}}
+  (void)__opencl_atomic_load(Ap, memory_order_relaxed, memory_scope_work_group);
+  (void)__opencl_atomic_load(Ap, memory_order_relaxed, memory_scope_device);
+  (void)__opencl_atomic_load(Ap, memory_order_relaxed, memory_scope_all_svm_devices);
+  (void)__opencl_atomic_load(Ap, memory_order_relaxed, memory_scope_sub_group);
+  (void)__opencl_atomic_load(Ap, memory_order_relaxed, scope);
+  (void)__opencl_atomic_load(Ap, memory_order_relaxed, 10);    //expected-error{{synchronization scope argument to atomic operation is invalid}}
+
+  // non-integer memory scope is casted to integer type.
+  float fscope;
+  (void)__opencl_atomic_load(Ap, memory_order_relaxed, 1.0f);
+  (void)__opencl_atomic_load(Ap, memory_order_relaxed, fscope);
+  struct S s;
+  (void)__opencl_atomic_load(Ap, memory_order_relaxed, s); //expected-error{{passing 'struct S' to parameter of incompatible type 'int'}}
+}
+
+void nullPointerWarning(atomic_int *Ap, int *p, int val) {
+  // The 'expected' pointer shouldn't be NULL.
+  (void)__opencl_atomic_compare_exchange_strong(Ap, (void *)0, val, memory_order_relaxed, memory_order_relaxed, memory_scope_work_group); // expected-warning {{null passed to a callee that requires a non-null argument}}
+}
diff --git a/test/SemaOpenCL/cl20-device-side-enqueue.cl b/test/SemaOpenCL/cl20-device-side-enqueue.cl
index 3f6527afeadc9..207fe7d340ab8 100644
--- a/test/SemaOpenCL/cl20-device-side-enqueue.cl
+++ b/test/SemaOpenCL/cl20-device-side-enqueue.cl
@@ -209,3 +209,35 @@ kernel void work_group_size_tests() {
   size = get_kernel_preferred_work_group_size_multiple(1);         // expected-error{{expected block argument}}
   size = get_kernel_preferred_work_group_size_multiple(block_A, 1); // expected-error{{too many arguments to function call, expected 1, have 2}}
 }
+
+#pragma OPENCL EXTENSION cl_khr_subgroups : enable
+
+kernel void foo(global int *buf)
+{
+  ndrange_t n;
+  buf[0] = get_kernel_max_sub_group_size_for_ndrange(n, ^(){});
+  buf[0] = get_kernel_max_sub_group_size_for_ndrange(0, ^(){}); // expected-error{{illegal call to 'get_kernel_max_sub_group_size_for_ndrange', expected 'ndrange_t' argument type}}
+  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)
+{
+  __private ndrange_t n;
+  buf[0] = get_kernel_sub_group_count_for_ndrange(n, ^(){});
+  buf[0] = get_kernel_sub_group_count_for_ndrange(0, ^(){}); // expected-error{{illegal call to 'get_kernel_sub_group_count_for_ndrange', expected 'ndrange_t' argument type}}
+  buf[0] = get_kernel_sub_group_count_for_ndrange(n, 1); // expected-error{{illegal call to 'get_kernel_sub_group_count_for_ndrange', expected block argument type}}
+}
+
+#pragma OPENCL EXTENSION cl_khr_subgroups : disable
+
+kernel void foo1(global 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)
+{
+  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 3e270e64c69d1..270345d86a617 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 -fsyntax-only -ferror-limit 100
+// RUN: %clang_cc1 %s -fblocks -verify -pedantic-errors -fsyntax-only -ferror-limit 100
 
 // Confirm CL2.0 Clang builtins are not available in earlier versions
 
diff --git a/test/SemaOpenCL/extension-version.cl b/test/SemaOpenCL/extension-version.cl
index 6100346f70729..e560276046513 100644
--- a/test/SemaOpenCL/extension-version.cl
+++ b/test/SemaOpenCL/extension-version.cl
@@ -273,3 +273,21 @@
 #endif
 #pragma OPENCL EXTENSION cl_amd_media_ops2: enable
 
+#if (__OPENCL_C_VERSION__ >= 120)
+#ifndef cl_intel_subgroups
+#error "Missing cl_intel_subgroups define"
+#endif
+#else
+// expected-warning@+2{{unsupported OpenCL extension 'cl_intel_subgroups' - ignoring}}
+#endif
+#pragma OPENCL EXTENSION cl_intel_subgroups : enable
+
+#if (__OPENCL_C_VERSION__ >= 120)
+#ifndef cl_intel_subgroups_short
+#error "Missing cl_intel_subgroups_short define"
+#endif
+#else
+// expected-warning@+2{{unsupported OpenCL extension 'cl_intel_subgroups_short' - ignoring}}
+#endif
+#pragma OPENCL EXTENSION cl_intel_subgroups_short : enable
+
diff --git a/test/SemaOpenCL/extern.cl b/test/SemaOpenCL/extern.cl
deleted file mode 100644
index 66efd7005836e..0000000000000
--- a/test/SemaOpenCL/extern.cl
+++ /dev/null
@@ -1,9 +0,0 @@
-// RUN: %clang_cc1 -x cl -cl-opt-disable -cl-std=CL1.2 -emit-llvm -ffake-address-space-map %s -o - -verify | FileCheck %s
-// expected-no-diagnostics
-
-// CHECK: @foo = external addrspace(2) constant float
-extern constant float foo;
-
-kernel void test(global float* buf) {
-  buf[0] += foo;
-}
diff --git a/test/SemaOpenCL/func.cl b/test/SemaOpenCL/func.cl
index dc5b44057b19f..83c3b4a6bcf90 100644
--- a/test/SemaOpenCL/func.cl
+++ b/test/SemaOpenCL/func.cl
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -triple spir-unknown-unknown
 
 // Variadic functions
 void vararg_f(int, ...);                    // expected-error {{invalid prototype, variadic arguments are not allowed in OpenCL}}
@@ -16,6 +16,9 @@ typedef struct s
 //Function pointer
 void foo(void*);
 
+// Expect no diagnostics for an empty parameter list.
+void bar();
+
 void bar()
 {
   // declaring a function pointer is an error
diff --git a/test/SemaOpenCL/invalid-block.cl b/test/SemaOpenCL/invalid-block.cl
index 89bf03264e13d..5d6dc380a37a1 100644
--- a/test/SemaOpenCL/invalid-block.cl
+++ b/test/SemaOpenCL/invalid-block.cl
@@ -81,3 +81,14 @@ kernel void f7() {
   };
   return;
 }
+
+// Taking address of a capture is not allowed
+int g;
+kernel void f8(int a1) {
+  int a2;
+  void (^bl)(void) = ^(void) {
+    &g; //expected-warning{{expression result unused}}
+    &a1; //expected-error{{taking address of a capture is not allowed}}
+    &a2; //expected-error{{taking address of a capture is not allowed}}
+  };
+}
diff --git a/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl b/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl
index 386c6b6c74561..619b359c7af98 100644
--- a/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl
+++ b/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl
@@ -1,4 +1,6 @@
-// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0 -cl-ext=+cl_khr_subgroups
+
+#pragma OPENCL EXTENSION cl_khr_subgroups : enable
 
 void test1(read_only pipe int p, global int* ptr){
   int tmp;
diff --git a/test/SemaOpenCL/invalid-pipes-cl2.0.cl b/test/SemaOpenCL/invalid-pipes-cl2.0.cl
index a50811650d2c7..463fd3d0dabc1 100644
--- a/test/SemaOpenCL/invalid-pipes-cl2.0.cl
+++ b/test/SemaOpenCL/invalid-pipes-cl2.0.cl
@@ -3,7 +3,7 @@
 global pipe int gp;            // expected-error {{type '__global read_only pipe int' can only be used as a function parameter in OpenCL}}
 global reserve_id_t rid;          // expected-error {{the '__global reserve_id_t' type cannot be used to declare a program scope variable}}
 
-extern pipe write_only int get_pipe(); // expected-error {{type '__global write_only pipe int ()' can only be used as a function parameter in OpenCL}}
+extern pipe write_only int get_pipe(); // expected-error {{type '__global write_only pipe int (void)' can only be used as a function parameter in OpenCL}}
 
 kernel void test_invalid_reserved_id(reserve_id_t ID) { // expected-error {{'reserve_id_t' cannot be used as the type of a kernel parameter}}
 }
diff --git a/test/SemaOpenCL/sampler_t.cl b/test/SemaOpenCL/sampler_t.cl
index 4d68dd20a17ec..5a1850b02cbd2 100644
--- a/test/SemaOpenCL/sampler_t.cl
+++ b/test/SemaOpenCL/sampler_t.cl
@@ -46,36 +46,11 @@ const constant sampler_t glb_smp11 = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_
 
 void kernel ker(sampler_t argsmp) {
   local sampler_t smp; // expected-error{{sampler type cannot be used with the __local and __global address space qualifiers}}
-  const sampler_t const_smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
-  const sampler_t const_smp2;
-  const sampler_t const_smp3 = const_smp;
-  const sampler_t const_smp4 = f();
   const sampler_t const_smp5 = 1.0f; // expected-error{{initializing 'const sampler_t' with an expression of incompatible type 'float'}}
   const sampler_t const_smp6 = 0x100000000LL; // expected-error{{sampler_t initialization requires 32-bit integer, not 'long long'}}
 
-  foo(glb_smp);
-  foo(glb_smp2);
-  foo(glb_smp3);
-  foo(glb_smp4);
-  foo(glb_smp5);
-  foo(glb_smp6);
-  foo(glb_smp7);
-  foo(glb_smp8);
-  foo(glb_smp9);
-  foo(smp);
-  foo(sampler_str.smp);
-  foo(const_smp);
-  foo(const_smp2);
-  foo(const_smp3);
-  foo(const_smp4);
-  foo(const_smp5);
-  foo(const_smp6);
-  foo(argsmp);
-  foo(5);
   foo(5.0f); // expected-error {{passing 'float' to parameter of incompatible type 'sampler_t'}}
-  sampler_t sa[] = {argsmp, const_smp}; // expected-error {{array of 'sampler_t' type is invalid in OpenCL}}
-  foo(sa[0]);
-  foo(bad());
+  sampler_t sa[] = {argsmp, glb_smp}; // expected-error {{array of 'sampler_t' type is invalid in OpenCL}}
 }
 
 void bad(sampler_t*); // expected-error{{pointer to type 'sampler_t' is invalid in OpenCL}}
diff --git a/test/SemaOpenCL/storageclass-cl20.cl b/test/SemaOpenCL/storageclass-cl20.cl
index b12676fe74377..581701d2a6a54 100644
--- a/test/SemaOpenCL/storageclass-cl20.cl
+++ b/test/SemaOpenCL/storageclass-cl20.cl
@@ -1,21 +1,41 @@
 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0
 
-static constant int G1 = 0;
 int G2 = 0;
 global int G3 = 0;
 local int G4 = 0;              // expected-error{{program scope variable must reside in global or constant address space}}
 
-void kernel foo() {
-  static int S1 = 5;
-  static global int S2 = 5;
-  static private int S3 = 5;   // expected-error{{static local variable must reside in global or constant address space}}
+static float g_implicit_static_var = 0;
+static constant float g_constant_static_var = 0;
+static global float g_global_static_var = 0;
+static local float g_local_static_var = 0;     // expected-error {{program scope variable must reside in global or constant address space}}
+static private float g_private_static_var = 0; // expected-error {{program scope variable must reside in global or constant address space}}
+static generic float g_generic_static_var = 0; // expected-error {{program scope variable must reside in global or constant address space}}
+
+extern float g_implicit_extern_var;
+extern constant float g_constant_extern_var;
+extern global float g_global_extern_var;
+extern local float g_local_extern_var;     // expected-error {{extern variable must reside in global or constant address space}}
+extern private float g_private_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
+extern generic float g_generic_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
 
+void kernel foo() {
   constant int L1 = 0;
   local int L2;
   global int L3;                              // expected-error{{function scope variable cannot be declared in global address space}}
   generic int L4;                             // expected-error{{automatic variable qualified with an invalid address space}}
   __attribute__((address_space(100))) int L5; // expected-error{{automatic variable qualified with an invalid address space}}
 
-  extern global int G5;
-  extern int G6; // expected-error{{extern variable must reside in global or constant address space}}
+  static float l_implicit_static_var = 0;
+  static constant float l_constant_static_var = 0;
+  static global float l_global_static_var = 0;
+  static local float l_local_static_var = 0;     // expected-error {{static local variable must reside in global or constant address space}}
+  static private float l_private_static_var = 0; // expected-error {{static local variable must reside in global or constant address space}}
+  static generic float l_generic_static_var = 0; // expected-error {{static local variable must reside in global or constant address space}}
+
+  extern float l_implicit_extern_var;
+  extern constant float l_constant_extern_var;
+  extern global float l_global_extern_var;
+  extern local float l_local_extern_var;     // expected-error {{extern variable must reside in global or constant address space}}
+  extern private float l_private_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
+  extern generic float l_generic_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
 }
diff --git a/test/SemaOpenCL/storageclass.cl b/test/SemaOpenCL/storageclass.cl
index 9a461068f2375..cbcb94509b320 100644
--- a/test/SemaOpenCL/storageclass.cl
+++ b/test/SemaOpenCL/storageclass.cl
@@ -5,6 +5,20 @@ constant int G2 = 0;
 int G3 = 0;        // expected-error{{program scope variable must reside in constant address space}}
 global int G4 = 0; // expected-error{{program scope variable must reside in constant address space}}
 
+static float g_implicit_static_var = 0; // expected-error {{program scope variable must reside in constant address space}}
+static constant float g_constant_static_var = 0;
+static global float g_global_static_var = 0;   // expected-error {{program scope variable must reside in constant address space}}
+static local float g_local_static_var = 0;     // expected-error {{program scope variable must reside in constant address space}}
+static private float g_private_static_var = 0; // expected-error {{program scope variable must reside in constant address space}}
+static generic float g_generic_static_var = 0; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{program scope variable must reside in constant address space}}
+
+extern float g_implicit_extern_var; // expected-error {{extern variable must reside in constant address space}}
+extern constant float g_constant_extern_var;
+extern global float g_global_extern_var;   // expected-error {{extern variable must reside in constant address space}}
+extern local float g_local_extern_var;     // expected-error {{extern variable must reside in constant address space}}
+extern private float g_private_extern_var; // expected-error {{extern variable must reside in constant address space}}
+extern generic float g_generic_extern_var; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{extern variable must reside in constant address space}}
+
 void kernel foo(int x) {
   // static is not allowed at local scope before CL2.0
   static int S1 = 5;          // expected-error{{variables in function scope cannot be declared static}}
@@ -45,10 +59,17 @@ void f() {
     __attribute__((address_space(100))) int L4; // expected-error{{automatic variable qualified with an invalid address space}}
   }
 
+  static float l_implicit_static_var = 0;          // expected-error {{variables in function scope cannot be declared static}}
+  static constant float l_constant_static_var = 0; // expected-error {{variables in function scope cannot be declared static}}
+  static global float l_global_static_var = 0;     // expected-error {{variables in function scope cannot be declared static}}
+  static local float l_local_static_var = 0;       // expected-error {{variables in function scope cannot be declared static}}
+  static private float l_private_static_var = 0;   // expected-error {{variables in function scope cannot be declared static}}
+  static generic float l_generic_static_var = 0;   // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{variables in function scope cannot be declared static}}
 
-  extern constant float L5;
-  extern local float L6; // expected-error{{extern variable must reside in constant address space}}
-
-  static int L7 = 0;     // expected-error{{variables in function scope cannot be declared static}}
-  static int L8;         // expected-error{{variables in function scope cannot be declared static}}
+  extern float l_implicit_extern_var; // expected-error {{extern variable must reside in constant address space}}
+  extern constant float l_constant_extern_var;
+  extern global float l_global_extern_var;   // expected-error {{extern variable must reside in constant address space}}
+  extern local float l_local_extern_var;     // expected-error {{extern variable must reside in constant address space}}
+  extern private float l_private_extern_var; // expected-error {{extern variable must reside in constant address space}}
+  extern generic float l_generic_extern_var; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{extern variable must reside in constant address space}}
 }
diff --git a/test/SemaOpenCL/to_addr_builtin.cl b/test/SemaOpenCL/to_addr_builtin.cl
index 56db4abed5790..70956f007a5fd 100644
--- a/test/SemaOpenCL/to_addr_builtin.cl
+++ b/test/SemaOpenCL/to_addr_builtin.cl
@@ -10,7 +10,7 @@ void test(void) {
 
   glob = to_global(glob, loc);
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
-  // expected-error@-2{{implicit declaration of function 'to_global' is invalid in OpenCL}}
+  // expected-warning@-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'}}
diff --git a/test/SemaOpenCL/vector_conv_invalid.cl b/test/SemaOpenCL/vector_conv_invalid.cl
index 90cec26a605a4..e32d84fd6a7e2 100644
--- a/test/SemaOpenCL/vector_conv_invalid.cl
+++ b/test/SemaOpenCL/vector_conv_invalid.cl
@@ -5,10 +5,18 @@ typedef int int4 __attribute((ext_vector_type(4)));
 typedef int int3 __attribute((ext_vector_type(3)));
 typedef unsigned uint3 __attribute((ext_vector_type(3)));
 
-void vector_conv_invalid() {
+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 e = (int4)u; // expected-error{{invalid conversion between ext-vector type 'int4' (vector of 4 'int' values) and 'uint4' (vector of 4 'unsigned int' values)}}
 
   uint3 u4 = (uint3)u; // expected-error{{invalid conversion between ext-vector type 'uint3' (vector of 3 'unsigned int' values) and 'uint4' (vector of 4 'unsigned int' values)}}
+
+  e = (const int4)i;
+  e = (constant int4)i;
+  e = (private int4)i;
+
+  private int4 *private_ptr = (const private int4 *)const_global_ptr; // expected-error{{casting 'const __global int4 *' to type 'const int4 *' changes address space of pointer}}
+  global int4 *global_ptr = const_global_ptr;                 // expected-warning {{initializing '__global int4 *' with an expression of type 'const __global int4 *' discards qualifiers}}
+  global_ptr = (global int4 *)const_global_ptr;
 }
diff --git a/test/SemaOpenCL/vector_swizzle_length.cl b/test/SemaOpenCL/vector_swizzle_length.cl
index 94e3f654d5d9c..4dbb5bd76ee0a 100644
--- a/test/SemaOpenCL/vector_swizzle_length.cl
+++ b/test/SemaOpenCL/vector_swizzle_length.cl
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// RUN: %clang_cc1 -x cl %s -verify -pedantic -fsyntax-only
 
 typedef float float8 __attribute__((ext_vector_type(8)));