9 files changed, 148 insertions, 0 deletions
diff --git a/clang/test/SemaOpenCL/address-spaces.cl b/clang/test/SemaOpenCL/address-spaces.cl
new file mode 100644
index 0000000..6ab10b3
--- /dev/null
+++ b/clang/test/SemaOpenCL/address-spaces.cl
@@ -0,0 +1,13 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+
+__constant int ci = 1;
+
+__kernel void foo(__global int *gip) {
+  __local int li;
+  __local int lj = 2; // expected-error {{'__local' variable cannot have an initializer}}
+
+  int *ip;
+  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}}
+}
diff --git a/clang/test/SemaOpenCL/cond.cl b/clang/test/SemaOpenCL/cond.cl
new file mode 100644
index 0000000..79dc82d
--- /dev/null
+++ b/clang/test/SemaOpenCL/cond.cl
@@ -0,0 +1,5 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+
+typedef __attribute__((ext_vector_type(4))) float float4;
+
+float4 foo(float4 a, float4 b, float4 c, float4 d) { return a < b ? c : d; }
diff --git a/clang/test/SemaOpenCL/extension-fp64.cl b/clang/test/SemaOpenCL/extension-fp64.cl
new file mode 100644
index 0000000..e0c2b1e
--- /dev/null
+++ b/clang/test/SemaOpenCL/extension-fp64.cl
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+
+void f1(double da) { // expected-error {{type 'double' requires cl_khr_fp64 extension}}
+  double d; // expected-error {{type 'double' requires cl_khr_fp64 extension}}
+  (void) 1.0; // expected-warning {{double precision constant requires cl_khr_fp64}}
+}
+
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+void f2(void) {
+  double d;
+  (void) 1.0;
+}
+
+#pragma OPENCL EXTENSION cl_khr_fp64 : disable
+
+void f3(void) {
+  double d; // expected-error {{type 'double' requires cl_khr_fp64 extension}}
+}
diff --git a/clang/test/SemaOpenCL/init.cl b/clang/test/SemaOpenCL/init.cl
new file mode 100644
index 0000000..b3ecfec
--- /dev/null
+++ b/clang/test/SemaOpenCL/init.cl
@@ -0,0 +1,15 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+
+typedef float float8 __attribute((ext_vector_type(8)));
+
+typedef float float32_t;
+typedef __attribute__(( __vector_size__(16) )) float32_t __neon_float32x4_t;
+typedef struct __simd128_float32_t {
+  __neon_float32x4_t val;
+} float32x4_t;
+
+float8 foo(float8 x) { 
+  float32x4_t lo;
+  float32x4_t hi;
+  return (float8) (lo.val, hi.val);
+}
diff --git a/clang/test/SemaOpenCL/vec_compare.cl b/clang/test/SemaOpenCL/vec_compare.cl
new file mode 100644
index 0000000..dd91aa5
--- /dev/null
+++ b/clang/test/SemaOpenCL/vec_compare.cl
@@ -0,0 +1,11 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+
+typedef __attribute__((ext_vector_type(2)))  unsigned int uint2;
+typedef __attribute__((ext_vector_type(2)))  int int2;
+
+void unsignedCompareOps()
+{
+  uint2 A, B;
+  int2 result = A != B;
+}
+
diff --git a/clang/test/SemaOpenCL/vec_step.cl b/clang/test/SemaOpenCL/vec_step.cl
new file mode 100644
index 0000000..d83ebf1
--- /dev/null
+++ b/clang/test/SemaOpenCL/vec_step.cl
@@ -0,0 +1,32 @@
+// RUN: %clang_cc1 -fsyntax-only -pedantic -verify %s
+
+typedef int int2 __attribute__((ext_vector_type(2)));
+typedef int int3 __attribute__((ext_vector_type(3)));
+typedef int int4 __attribute__((ext_vector_type(4)));
+typedef int int8 __attribute__((ext_vector_type(8)));
+typedef int int16 __attribute__((ext_vector_type(16)));
+
+void foo(int3 arg1, int8 arg2) {
+  int4 auto1;
+  int16 *auto2;
+  int auto3;
+  int2 auto4;
+  struct S *incomplete1;
+
+  int res1[vec_step(arg1) == 4 ? 1 : -1];
+  int res2[vec_step(arg2) == 8 ? 1 : -1];
+  int res3[vec_step(auto1) == 4 ? 1 : -1];
+  int res4[vec_step(*auto2) == 16 ? 1 : -1];
+  int res5[vec_step(auto3) == 1 ? 1 : -1];
+  int res6[vec_step(auto4) == 2 ? 1 : -1];
+  int res7[vec_step(int2) == 2 ? 1 : -1];
+  int res8[vec_step(int3) == 4 ? 1 : -1];
+  int res9[vec_step(int4) == 4 ? 1 : -1];
+  int res10[vec_step(int8) == 8 ? 1 : -1];
+  int res11[vec_step(int16) == 16 ? 1 : -1];
+  int res12[vec_step(void) == 1 ? 1 : -1];
+
+  int res13 = vec_step(*incomplete1); // expected-error {{'vec_step' requires built-in scalar or vector type, 'struct S' invalid}}
+  int res14 = vec_step(int16*); // expected-error {{'vec_step' requires built-in scalar or vector type, 'int16 *' invalid}}
+  int res15 = vec_step(void(void)); // expected-error {{'vec_step' requires built-in scalar or vector type, 'void (void)' invalid}}
+}
diff --git a/clang/test/SemaOpenCL/vector_conv_invalid.cl b/clang/test/SemaOpenCL/vector_conv_invalid.cl
new file mode 100644
index 0000000..e6ef5a4
--- /dev/null
+++ b/clang/test/SemaOpenCL/vector_conv_invalid.cl
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -verify %s
+
+typedef unsigned int uint4 __attribute((ext_vector_type(4)));
+typedef int int4 __attribute((ext_vector_type(4)));
+typedef int int3 __attribute((ext_vector_type(3)));
+typedef unsigned uint3 __attribute((ext_vector_type(3)));
+
+void vector_conv_invalid() {
+  uint4 u = (uint4)(1);
+  int4 i = u; // expected-error{{initializing 'int4' with an expression of incompatible type 'uint4'}}
+  int4 e = (int4)u; // expected-error{{invalid conversion between ext-vector type 'int4' and 'uint4'}}
+
+  uint3 u4 = (uint3)u; // expected-error{{invalid conversion between ext-vector type 'uint3' and 'uint4'}}
+}
diff --git a/clang/test/SemaOpenCL/vector_literals_const.cl b/clang/test/SemaOpenCL/vector_literals_const.cl
new file mode 100644
index 0000000..e761816
--- /dev/null
+++ b/clang/test/SemaOpenCL/vector_literals_const.cl
@@ -0,0 +1,26 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+
+typedef int int2 __attribute((ext_vector_type(2)));
+typedef int int3 __attribute((ext_vector_type(3)));
+typedef int int4 __attribute((ext_vector_type(4)));
+
+__constant int4 i_1_1_1_1 = (int4)(1,2,3,4);
+__constant int4 i_2_1_1 = (int4)((int2)(1,2),3,4);
+__constant int4 i_1_2_1 = (int4)(1,(int2)(2,3),4);
+__constant int4 i_1_1_2 = (int4)(1,2,(int2)(3,4));
+__constant int4 i_2_2 = (int4)((int2)(1,2),(int2)(3,4));
+__constant int4 i_3_1 = (int4)((int3)(1,2,3),4);
+__constant int4 i_1_3 = (int4)(1,(int3)(2,3,4));
+
+typedef float float2 __attribute((ext_vector_type(2)));
+typedef float float3 __attribute((ext_vector_type(3)));
+typedef float float4 __attribute((ext_vector_type(4)));
+
+__constant float4 f_1_1_1_1 = (float4)(1,2,3,4);
+__constant float4 f_2_1_1 = (float4)((float2)(1,2),3,4);
+__constant float4 f_1_2_1 = (float4)(1,(float2)(2,3),4);
+__constant float4 f_1_1_2 = (float4)(1,2,(float2)(3,4));
+__constant float4 f_2_2 = (float4)((float2)(1,2),(float2)(3,4));
+__constant float4 f_3_1 = (float4)((float3)(1,2,3),4);
+__constant float4 f_1_3 = (float4)(1,(float3)(2,3,4));
+
diff --git a/clang/test/SemaOpenCL/vector_literals_invalid.cl b/clang/test/SemaOpenCL/vector_literals_invalid.cl
new file mode 100644
index 0000000..e4e23cd
--- /dev/null
+++ b/clang/test/SemaOpenCL/vector_literals_invalid.cl
@@ -0,0 +1,13 @@
+// RUN: %clang_cc1 -verify %s 
+
+typedef __attribute__(( ext_vector_type(4) ))  float float4;
+typedef __attribute__(( ext_vector_type(4) ))  int int4;
+typedef __attribute__(( ext_vector_type(8) ))  int int8;
+
+void vector_literals_invalid()
+{
+  int4 a = (int4)(1,2,3); // expected-error{{too few elements}}
+  int4 b = (int4)(1,2,3,4,5); // expected-error{{excess elements in vector}}
+  ((float4)(1.0f))++; // expected-error{{cannot increment value of type 'float4'}}
+  int8 d = (int8)(a,(float4)(1)); // expected-error{{initializing 'int' with an expression of incompatible type 'float4'}}
+}