vendor/clang/clang-release_39-r276489

30 files changed, 1116 insertions, 205 deletions

diff --git a/test/CodeGenOpenCL/address-spaces-conversions.cl b/test/CodeGenOpenCL/address-spaces-conversions.cl
index bc80f47f5f42..c947db41e070 100644
--- a/test/CodeGenOpenCL/address-spaces-conversions.cl
+++ b/test/CodeGenOpenCL/address-spaces-conversions.cl
@@ -1,22 +1,89 @@
 // RUN: %clang_cc1 %s -triple x86_64-unknown-linux-gnu -O0 -ffake-address-space-map -cl-std=CL2.0 -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple x86_64-unknown-linux-gnu -O0 -cl-std=CL2.0 -emit-llvm -o - | FileCheck --check-prefix=CHECK-NOFAKE %s
+// When -ffake-address-space-map is not used, all addr space mapped to 0 for x86_64.
 
 // test that we generate address space casts everywhere we need conversions of
 // pointers to different address spaces
 
+// CHECK: define void @test
 void test(global int *arg_glob, generic int *arg_gen) {
   int var_priv;
   arg_gen = arg_glob; // implicit cast global -> generic
   // CHECK: %{{[0-9]+}} = addrspacecast i32 addrspace(1)* %{{[0-9]+}} to i32 addrspace(4)*
+  // CHECK-NOFAKE-NOT: addrspacecast
+
   arg_gen = &var_priv; // implicit cast with obtaining adr, private -> generic
   // CHECK: %{{[0-9]+}} = addrspacecast i32* %var_priv to i32 addrspace(4)*
+  // CHECK-NOFAKE-NOT: addrspacecast
+
   arg_glob = (global int *)arg_gen; // explicit cast
   // CHECK: %{{[0-9]+}} = addrspacecast i32 addrspace(4)* %{{[0-9]+}} to i32 addrspace(1)*
+  // CHECK-NOFAKE-NOT: addrspacecast
+
   global int *var_glob =
       (global int *)arg_glob; // explicit cast in the same address space
   // CHECK-NOT: %{{[0-9]+}} = addrspacecast i32 addrspace(1)* %{{[0-9]+}} to i32 addrspace(1)*
+  // CHECK-NOFAKE-NOT: addrspacecast
+
   var_priv = arg_gen - arg_glob; // arithmetic operation
   // CHECK: %{{.*}} = ptrtoint i32 addrspace(4)* %{{.*}} to i64
   // CHECK: %{{.*}} = ptrtoint i32 addrspace(1)* %{{.*}} to i64
+  // CHECK-NOFAKE: %{{.*}} = ptrtoint i32* %{{.*}} to i64
+  // CHECK-NOFAKE: %{{.*}} = ptrtoint i32* %{{.*}} to i64
+
   var_priv = arg_gen > arg_glob; // comparison
   // CHECK: %{{[0-9]+}} = addrspacecast i32 addrspace(1)* %{{[0-9]+}} to i32 addrspace(4)*
+
+  generic void *var_gen_v = arg_glob;
+  // CHECK: addrspacecast
+  // CHECK-NOT: bitcast
+  // CHECK-NOFAKE: bitcast
+  // CHECK-NOFAKE-NOT: addrspacecast
+}
+
+// Test ternary operator.
+// CHECK: define void @test_ternary
+void test_ternary(void) {
+  global int *var_glob;
+  generic int *var_gen;
+  generic int *var_gen2;
+  generic float *var_gen_f;
+  generic void *var_gen_v;
+
+  var_gen = var_gen ? var_gen : var_gen2; // operands of the same addr spaces and the same type
+  // CHECK: icmp
+  // CHECK-NOT: addrspacecast
+  // CHECK-NOT: bitcast
+  // CHECK: phi
+  // CHECK: store i32 addrspace(4)* %{{.+}}, i32 addrspace(4)** %{{.+}}
+
+  var_gen = var_gen ? var_gen : var_glob; // operands of overlapping addr spaces and the same type
+  // CHECK: icmp
+  // CHECK-NOT: bitcast
+  // CHECK: %{{.+}} = addrspacecast i32 addrspace(1)* %{{.+}} to i32 addrspace(4)*
+  // CHECK: phi
+  // CHECK: store
+
+  typedef int int_t;
+  global int_t *var_glob_typedef;
+  var_gen = var_gen ? var_gen : var_glob_typedef; // operands of overlapping addr spaces and equivalent types
+  // CHECK: icmp
+  // CHECK-NOT: bitcast
+  // CHECK: %{{.+}} = addrspacecast i32 addrspace(1)* %{{.+}} to i32 addrspace(4)*
+  // CHECK: phi
+  // CHECK: store
+ 
+  var_gen_v = var_gen ? var_gen : var_gen_f; // operands of the same addr space and different types
+  // CHECK: icmp
+  // CHECK: %{{.+}} = bitcast i32 addrspace(4)* %{{.+}} to i8 addrspace(4)*
+  // CHECK: %{{.+}} = bitcast float addrspace(4)* %{{.+}} to i8 addrspace(4)*
+  // CHECK: phi
+  // CHECK: store
+
+  var_gen_v = var_gen ? var_glob : var_gen_f; // operands of overlapping addr spaces and different types
+  // CHECK: icmp
+  // CHECK: %{{.+}} = addrspacecast i32 addrspace(1)* %{{.+}} to i8 addrspace(4)*
+  // CHECK: %{{.+}} = bitcast float addrspace(4)* %{{.+}} to i8 addrspace(4)*
+  // CHECK: phi
+  // CHECK: store
 }
diff --git a/test/CodeGenOpenCL/amdgcn-flat-scratch-name.cl b/test/CodeGenOpenCL/amdgcn-flat-scratch-name.cl
new file mode 100644
index 000000000000..3a98e9099f52
--- /dev/null
+++ b/test/CodeGenOpenCL/amdgcn-flat-scratch-name.cl
@@ -0,0 +1,15 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
+
+// CHECK-LABEL: @use_flat_scratch_name
+kernel void use_flat_scratch_name()
+{
+// CHECK: tail call void asm sideeffect "s_mov_b64 flat_scratch, 0", "~{flat_scratch}"()
+  __asm__ volatile("s_mov_b64 flat_scratch, 0" : : : "flat_scratch");
+
+// CHECK: tail call void asm sideeffect "s_mov_b32 flat_scratch_lo, 0", "~{flat_scratch_lo}"()
+  __asm__ volatile("s_mov_b32 flat_scratch_lo, 0" : : : "flat_scratch_lo");
+
+// CHECK: tail call void asm sideeffect "s_mov_b32 flat_scratch_hi, 0", "~{flat_scratch_hi}"()
+  __asm__ volatile("s_mov_b32 flat_scratch_hi, 0" : : : "flat_scratch_hi");
+}
diff --git a/test/CodeGenOpenCL/amdgpu-call-kernel.cl b/test/CodeGenOpenCL/amdgpu-call-kernel.cl
new file mode 100755
index 000000000000..005793916c68
--- /dev/null
+++ b/test/CodeGenOpenCL/amdgpu-call-kernel.cl
@@ -0,0 +1,14 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
+// CHECK: define amdgpu_kernel void @test_call_kernel(i32 addrspace(1)* nocapture %out)
+// CHECK: store i32 4, i32 addrspace(1)* %out, align 4
+
+kernel void test_kernel(global int *out)
+{
+  out[0] = 4;
+}
+
+__kernel void test_call_kernel(__global int *out)
+{
+  test_kernel(out);
+}
diff --git a/test/CodeGenOpenCL/amdgpu-calling-conv.cl b/test/CodeGenOpenCL/amdgpu-calling-conv.cl
new file mode 100644
index 000000000000..7da9d7f4d49a
--- /dev/null
+++ b/test/CodeGenOpenCL/amdgpu-calling-conv.cl
@@ -0,0 +1,12 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
+
+// CHECK: define amdgpu_kernel void @calling_conv_amdgpu_kernel()
+kernel void calling_conv_amdgpu_kernel()
+{
+}
+
+// CHECK: define void @calling_conv_none()
+void calling_conv_none()
+{
+}
diff --git a/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl b/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl
index 35bdcead3128..589d00d1eaa4 100644
--- a/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl
+++ b/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl
@@ -5,23 +5,23 @@
 
 __attribute__((amdgpu_num_vgpr(64))) // expected-no-diagnostics
 kernel void test_num_vgpr64() {
-// CHECK: define void @test_num_vgpr64() [[ATTR_VGPR64:#[0-9]+]]
+// CHECK: define amdgpu_kernel void @test_num_vgpr64() [[ATTR_VGPR64:#[0-9]+]]
 }
 
 __attribute__((amdgpu_num_sgpr(32))) // expected-no-diagnostics
 kernel void test_num_sgpr32() {
-// CHECK: define void @test_num_sgpr32() [[ATTR_SGPR32:#[0-9]+]]
+// CHECK: define amdgpu_kernel void @test_num_sgpr32() [[ATTR_SGPR32:#[0-9]+]]
 }
 
 __attribute__((amdgpu_num_vgpr(64), amdgpu_num_sgpr(32))) // expected-no-diagnostics
 kernel void test_num_vgpr64_sgpr32() {
-// CHECK: define void @test_num_vgpr64_sgpr32() [[ATTR_VGPR64_SGPR32:#[0-9]+]]
+// CHECK: define amdgpu_kernel void @test_num_vgpr64_sgpr32() [[ATTR_VGPR64_SGPR32:#[0-9]+]]
 
 }
 
 __attribute__((amdgpu_num_sgpr(20), amdgpu_num_vgpr(40))) // expected-no-diagnostics
 kernel void test_num_sgpr20_vgpr40() {
-// CHECK: define void @test_num_sgpr20_vgpr40() [[ATTR_SGPR20_VGPR40:#[0-9]+]]
+// CHECK: define amdgpu_kernel void @test_num_sgpr20_vgpr40() [[ATTR_SGPR20_VGPR40:#[0-9]+]]
 }
 
 __attribute__((amdgpu_num_vgpr(0))) // expected-no-diagnostics
@@ -40,8 +40,8 @@ kernel void test_num_vgpr0_sgpr0() {
 // X86-NOT: "amdgpu_num_vgpr"
 // X86-NOT: "amdgpu_num_sgpr"
 
-// CHECK-DAG-NOT: "amdgpu_num_vgpr"="0"
-// CHECK-DAG-NOT: "amdgpu_num_sgpr"="0"
+// CHECK-NOT: "amdgpu_num_vgpr"="0"
+// CHECK-NOT: "amdgpu_num_sgpr"="0"
 // CHECK-DAG: attributes [[ATTR_VGPR64]] = { nounwind "amdgpu_num_vgpr"="64"
 // CHECK-DAG: attributes [[ATTR_SGPR32]] = { nounwind "amdgpu_num_sgpr"="32"
 // CHECK-DAG: attributes [[ATTR_VGPR64_SGPR32]] = { nounwind "amdgpu_num_sgpr"="32" "amdgpu_num_vgpr"="64"
diff --git a/test/CodeGenOpenCL/as_type.cl b/test/CodeGenOpenCL/as_type.cl
new file mode 100644
index 000000000000..7fc3b02bdc75
--- /dev/null
+++ b/test/CodeGenOpenCL/as_type.cl
@@ -0,0 +1,68 @@
+// RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -o - | FileCheck %s
+
+typedef __attribute__(( ext_vector_type(3) )) char char3;
+typedef __attribute__(( ext_vector_type(4) )) char char4;
+typedef __attribute__(( ext_vector_type(16) )) char char16;
+typedef __attribute__(( ext_vector_type(3) )) int int3;
+
+//CHECK: define spir_func <3 x i8> @f1(<4 x i8> %[[x:.*]])
+//CHECK: %[[astype:.*]] = shufflevector <4 x i8> %[[x]], <4 x i8> undef, <3 x i32> <i32 0, i32 1, i32 2>
+//CHECK: ret <3 x i8> %[[astype]]
+char3 f1(char4 x) {
+  return  __builtin_astype(x, char3);
+}
+
+//CHECK: define spir_func <4 x i8> @f2(<3 x i8> %[[x:.*]])
+//CHECK: %[[astype:.*]] = shufflevector <3 x i8> %[[x]], <3 x i8> undef, <4 x i32> <i32 0, i32 1, i32 2, i32 undef>
+//CHECK: ret <4 x i8> %[[astype]]
+char4 f2(char3 x) {
+  return __builtin_astype(x, char4);
+}
+
+//CHECK: define spir_func <3 x i8> @f3(i32 %[[x:.*]])
+//CHECK: %[[cast:.*]] = bitcast i32 %[[x]] to <4 x i8>
+//CHECK: %[[astype:.*]] = shufflevector <4 x i8> %[[cast]], <4 x i8> undef, <3 x i32> <i32 0, i32 1, i32 2>
+//CHECK: ret <3 x i8> %[[astype]]
+char3 f3(int x) {
+  return __builtin_astype(x, char3);
+}
+
+//CHECK: define spir_func <4 x i8> @f4(i32 %[[x:.*]])
+//CHECK: %[[astype:.*]] = bitcast i32 %[[x]] to <4 x i8>
+//CHECK-NOT: shufflevector
+//CHECK: ret <4 x i8> %[[astype]]
+char4 f4(int x) {
+  return __builtin_astype(x, char4);
+}
+
+//CHECK: define spir_func i32 @f5(<3 x i8> %[[x:.*]])
+//CHECK: %[[shuffle:.*]] = shufflevector <3 x i8> %[[x]], <3 x i8> undef, <4 x i32> <i32 0, i32 1, i32 2, i32 undef>
+//CHECK: %[[astype:.*]] = bitcast <4 x i8> %[[shuffle]] to i32
+//CHECK: ret i32 %[[astype]]
+int f5(char3 x) {
+  return __builtin_astype(x, int);
+}
+
+//CHECK: define spir_func i32 @f6(<4 x i8> %[[x:.*]])
+//CHECK: %[[astype:.*]] = bitcast <4 x i8> %[[x]] to i32
+//CHECK-NOT: shufflevector
+//CHECK: ret i32 %[[astype]]
+int f6(char4 x) {
+  return __builtin_astype(x, int);
+}
+
+//CHECK: define spir_func <3 x i8> @f7(<3 x i8> %[[x:.*]])
+//CHECK-NOT: bitcast
+//CHECK-NOT: shufflevector
+//CHECK: ret <3 x i8> %[[x]]
+char3 f7(char3 x) {
+  return __builtin_astype(x, char3);
+}
+
+//CHECK: define spir_func <3 x i32> @f8(<16 x i8> %[[x:.*]])
+//CHECK: %[[cast:.*]] = bitcast <16 x i8> %[[x]] to <4 x i32>
+//CHECK: %[[astype:.*]] = shufflevector <4 x i32> %[[cast]], <4 x i32> undef, <3 x i32> <i32 0, i32 1, i32 2>
+//CHECK: ret <3 x i32> %[[astype]]
+int3 f8(char16 x) {
+  return __builtin_astype(x, int3);
+}
diff --git a/test/CodeGenOpenCL/builtins-amdgcn-error.cl b/test/CodeGenOpenCL/builtins-amdgcn-error.cl
new file mode 100644
index 000000000000..89c3e490ecdd
--- /dev/null
+++ b/test/CodeGenOpenCL/builtins-amdgcn-error.cl
@@ -0,0 +1,18 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-amdhsa -target-cpu tahiti -verify -S -o - %s
+
+// FIXME: We only get one error if the functions are the other order in the
+// file.
+
+typedef unsigned long ulong;
+
+ulong test_s_memrealtime()
+{
+  return __builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
+}
+
+void test_s_sleep(int x)
+{
+  __builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}}
+}
+
diff --git a/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
new file mode 100644
index 000000000000..cda87a8e1ed1
--- /dev/null
+++ b/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -0,0 +1,12 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -o - %s | FileCheck %s
+
+typedef unsigned long ulong;
+
+
+// CHECK-LABEL: @test_s_memrealtime
+// CHECK: call i64 @llvm.amdgcn.s.memrealtime()
+void test_s_memrealtime(global ulong* out)
+{
+  *out = __builtin_amdgcn_s_memrealtime();
+}
diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl
new file mode 100644
index 000000000000..6cac2a44bc79
--- /dev/null
+++ b/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -0,0 +1,316 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
+
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+typedef unsigned long ulong;
+
+// CHECK-LABEL: @test_div_scale_f64
+// CHECK: call { double, i1 } @llvm.amdgcn.div.scale.f64(double %a, double %b, i1 true)
+// CHECK-DAG: [[FLAG:%.+]] = extractvalue { double, i1 } %{{.+}}, 1
+// CHECK-DAG: [[VAL:%.+]] = extractvalue { double, i1 } %{{.+}}, 0
+// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i32
+// CHECK: store i32 [[FLAGEXT]]
+void test_div_scale_f64(global double* out, global int* flagout, double a, double b)
+{
+  bool flag;
+  *out = __builtin_amdgcn_div_scale(a, b, true, &flag);
+  *flagout = flag;
+}
+
+// CHECK-LABEL: @test_div_scale_f32
+// CHECK: call { float, i1 } @llvm.amdgcn.div.scale.f32(float %a, float %b, i1 true)
+// CHECK-DAG: [[FLAG:%.+]] = extractvalue { float, i1 } %{{.+}}, 1
+// CHECK-DAG: [[VAL:%.+]] = extractvalue { float, i1 } %{{.+}}, 0
+// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i32
+// CHECK: store i32 [[FLAGEXT]]
+void test_div_scale_f32(global float* out, global int* flagout, float a, float b)
+{
+  bool flag;
+  *out = __builtin_amdgcn_div_scalef(a, b, true, &flag);
+  *flagout = flag;
+}
+
+// CHECK-LABEL: @test_div_fmas_f32
+// CHECK: call float @llvm.amdgcn.div.fmas.f32
+void test_div_fmas_f32(global float* out, float a, float b, float c, int d)
+{
+  *out = __builtin_amdgcn_div_fmasf(a, b, c, d);
+}
+
+// CHECK-LABEL: @test_div_fmas_f64
+// CHECK: call double @llvm.amdgcn.div.fmas.f64
+void test_div_fmas_f64(global double* out, double a, double b, double c, int d)
+{
+  *out = __builtin_amdgcn_div_fmas(a, b, c, d);
+}
+
+// CHECK-LABEL: @test_div_fixup_f32
+// CHECK: call float @llvm.amdgcn.div.fixup.f32
+void test_div_fixup_f32(global float* out, float a, float b, float c)
+{
+  *out = __builtin_amdgcn_div_fixupf(a, b, c);
+}
+
+// CHECK-LABEL: @test_div_fixup_f64
+// CHECK: call double @llvm.amdgcn.div.fixup.f64
+void test_div_fixup_f64(global double* out, double a, double b, double c)
+{
+  *out = __builtin_amdgcn_div_fixup(a, b, c);
+}
+
+// CHECK-LABEL: @test_trig_preop_f32
+// CHECK: call float @llvm.amdgcn.trig.preop.f32
+void test_trig_preop_f32(global float* out, float a, int b)
+{
+  *out = __builtin_amdgcn_trig_preopf(a, b);
+}
+
+// CHECK-LABEL: @test_trig_preop_f64
+// CHECK: call double @llvm.amdgcn.trig.preop.f64
+void test_trig_preop_f64(global double* out, double a, int b)
+{
+  *out = __builtin_amdgcn_trig_preop(a, b);
+}
+
+// CHECK-LABEL: @test_rcp_f32
+// CHECK: call float @llvm.amdgcn.rcp.f32
+void test_rcp_f32(global float* out, float a)
+{
+  *out = __builtin_amdgcn_rcpf(a);
+}
+
+// CHECK-LABEL: @test_rcp_f64
+// CHECK: call double @llvm.amdgcn.rcp.f64
+void test_rcp_f64(global double* out, double a)
+{
+  *out = __builtin_amdgcn_rcp(a);
+}
+
+// CHECK-LABEL: @test_rsq_f32
+// CHECK: call float @llvm.amdgcn.rsq.f32
+void test_rsq_f32(global float* out, float a)
+{
+  *out = __builtin_amdgcn_rsqf(a);
+}
+
+// CHECK-LABEL: @test_rsq_f64
+// CHECK: call double @llvm.amdgcn.rsq.f64
+void test_rsq_f64(global double* out, double a)
+{
+  *out = __builtin_amdgcn_rsq(a);
+}
+
+// CHECK-LABEL: @test_rsq_clamp_f32
+// CHECK: call float @llvm.amdgcn.rsq.clamp.f32
+void test_rsq_clamp_f32(global float* out, float a)
+{
+  *out = __builtin_amdgcn_rsq_clampf(a);
+}
+
+// CHECK-LABEL: @test_rsq_clamp_f64
+// CHECK: call double @llvm.amdgcn.rsq.clamp.f64
+void test_rsq_clamp_f64(global double* out, double a)
+{
+  *out = __builtin_amdgcn_rsq_clamp(a);
+}
+
+// CHECK-LABEL: @test_sin_f32
+// CHECK: call float @llvm.amdgcn.sin.f32
+void test_sin_f32(global float* out, float a)
+{
+  *out = __builtin_amdgcn_sinf(a);
+}
+
+// CHECK-LABEL: @test_cos_f32
+// CHECK: call float @llvm.amdgcn.cos.f32
+void test_cos_f32(global float* out, float a)
+{
+  *out = __builtin_amdgcn_cosf(a);
+}
+
+// CHECK-LABEL: @test_log_clamp_f32
+// CHECK: call float @llvm.amdgcn.log.clamp.f32
+void test_log_clamp_f32(global float* out, float a)
+{
+  *out = __builtin_amdgcn_log_clampf(a);
+}
+
+// CHECK-LABEL: @test_ldexp_f32
+// CHECK: call float @llvm.amdgcn.ldexp.f32
+void test_ldexp_f32(global float* out, float a, int b)
+{
+  *out = __builtin_amdgcn_ldexpf(a, b);
+}
+
+// CHECK-LABEL: @test_ldexp_f64
+// CHECK: call double @llvm.amdgcn.ldexp.f64
+void test_ldexp_f64(global double* out, double a, int b)
+{
+  *out = __builtin_amdgcn_ldexp(a, b);
+}
+
+// CHECK-LABEL: @test_frexp_mant_f32
+// CHECK: call float @llvm.amdgcn.frexp.mant.f32
+void test_frexp_mant_f32(global float* out, float a)
+{
+  *out = __builtin_amdgcn_frexp_mantf(a);
+}
+
+// CHECK-LABEL: @test_frexp_mant_f64
+// CHECK: call double @llvm.amdgcn.frexp.mant.f64
+void test_frexp_mant_f64(global double* out, double a)
+{
+  *out = __builtin_amdgcn_frexp_mant(a);
+}
+
+// CHECK-LABEL: @test_frexp_exp_f32
+// CHECK: call i32 @llvm.amdgcn.frexp.exp.f32
+void test_frexp_exp_f32(global int* out, float a)
+{
+  *out = __builtin_amdgcn_frexp_expf(a);
+}
+
+// CHECK-LABEL: @test_frexp_exp_f64
+// CHECK: call i32 @llvm.amdgcn.frexp.exp.f64
+void test_frexp_exp_f64(global int* out, double a)
+{
+  *out = __builtin_amdgcn_frexp_exp(a);
+}
+
+// CHECK-LABEL: @test_fract_f32
+// CHECK: call float @llvm.amdgcn.fract.f32
+void test_fract_f32(global int* out, float a)
+{
+  *out = __builtin_amdgcn_fractf(a);
+}
+
+// CHECK-LABEL: @test_fract_f64
+// CHECK: call double @llvm.amdgcn.fract.f64
+void test_fract_f64(global int* out, double a)
+{
+  *out = __builtin_amdgcn_fract(a);
+}
+
+// CHECK-LABEL: @test_lerp
+// CHECK: call i32 @llvm.amdgcn.lerp
+void test_lerp(global int* out, int a, int b, int c)
+{
+  *out = __builtin_amdgcn_lerp(a, b, c);
+}
+
+// CHECK-LABEL: @test_class_f32
+// CHECK: call i1 @llvm.amdgcn.class.f32
+void test_class_f32(global float* out, float a, int b)
+{
+  *out = __builtin_amdgcn_classf(a, b);
+}
+
+// CHECK-LABEL: @test_class_f64
+// CHECK: call i1 @llvm.amdgcn.class.f64
+void test_class_f64(global double* out, double a, int b)
+{
+  *out = __builtin_amdgcn_class(a, b);
+}
+
+// CHECK-LABEL: @test_s_barrier
+// CHECK: call void @llvm.amdgcn.s.barrier(
+void test_s_barrier()
+{
+  __builtin_amdgcn_s_barrier();
+}
+
+// CHECK-LABEL: @test_s_memtime
+// CHECK: call i64 @llvm.amdgcn.s.memtime()
+void test_s_memtime(global ulong* out)
+{
+  *out = __builtin_amdgcn_s_memtime();
+}
+
+// CHECK-LABEL: @test_s_sleep
+// CHECK: call void @llvm.amdgcn.s.sleep(i32 1)
+// CHECK: call void @llvm.amdgcn.s.sleep(i32 15)
+void test_s_sleep()
+{
+  __builtin_amdgcn_s_sleep(1);
+  __builtin_amdgcn_s_sleep(15);
+}
+
+// CHECK-LABEL: @test_cubeid(
+// CHECK: call float @llvm.amdgcn.cubeid(float %a, float %b, float %c)
+void test_cubeid(global float* out, float a, float b, float c) {
+  *out = __builtin_amdgcn_cubeid(a, b, c);
+}
+
+// CHECK-LABEL: @test_cubesc(
+// CHECK: call float @llvm.amdgcn.cubesc(float %a, float %b, float %c)
+void test_cubesc(global float* out, float a, float b, float c) {
+  *out = __builtin_amdgcn_cubesc(a, b, c);
+}
+
+// CHECK-LABEL: @test_cubetc(
+// CHECK: call float @llvm.amdgcn.cubetc(float %a, float %b, float %c)
+void test_cubetc(global float* out, float a, float b, float c) {
+  *out = __builtin_amdgcn_cubetc(a, b, c);
+}
+
+// CHECK-LABEL: @test_cubema(
+// CHECK: call float @llvm.amdgcn.cubema(float %a, float %b, float %c)
+void test_cubema(global float* out, float a, float b, float c) {
+  *out = __builtin_amdgcn_cubema(a, b, c);
+}
+
+// CHECK-LABEL: @test_read_exec(
+// CHECK: call i64 @llvm.read_register.i64(metadata ![[EXEC:[0-9]+]]) #[[READ_EXEC_ATTRS:[0-9]+]]
+void test_read_exec(global ulong* out) {
+  *out = __builtin_amdgcn_read_exec();
+}
+
+// CHECK: declare i64 @llvm.read_register.i64(metadata) #[[NOUNWIND_READONLY:[0-9]+]]
+
+// CHECK-LABEL: @test_kernarg_segment_ptr
+// CHECK: call i8 addrspace(2)* @llvm.amdgcn.kernarg.segment.ptr()
+void test_kernarg_segment_ptr(__attribute__((address_space(2))) unsigned char ** out)
+{
+  *out = __builtin_amdgcn_kernarg_segment_ptr();
+}
+
+// CHECK-LABEL: @test_implicitarg_ptr
+// CHECK: call i8 addrspace(2)* @llvm.amdgcn.implicitarg.ptr()
+void test_implicitarg_ptr(__attribute__((address_space(2))) unsigned char ** out)
+{
+  *out = __builtin_amdgcn_implicitarg_ptr();
+}
+
+// CHECK-LABEL: @test_get_group_id(
+// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.x()
+// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.y()
+// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.z()
+void test_get_group_id(int d, global int *out)
+{
+	switch (d) {
+	case 0: *out = __builtin_amdgcn_workgroup_id_x(); break;
+	case 1: *out = __builtin_amdgcn_workgroup_id_y(); break;
+	case 2: *out = __builtin_amdgcn_workgroup_id_z(); break;
+	default: *out = 0;
+	}
+}
+
+// CHECK-LABEL: @test_get_local_id(
+// CHECK: tail call i32 @llvm.amdgcn.workitem.id.x(), !range [[WI_RANGE:![0-9]*]]
+// CHECK: tail call i32 @llvm.amdgcn.workitem.id.y(), !range [[WI_RANGE]]
+// CHECK: tail call i32 @llvm.amdgcn.workitem.id.z(), !range [[WI_RANGE]]
+void test_get_local_id(int d, global int *out)
+{
+	switch (d) {
+	case 0: *out = __builtin_amdgcn_workitem_id_x(); break;
+	case 1: *out = __builtin_amdgcn_workitem_id_y(); break;
+	case 2: *out = __builtin_amdgcn_workitem_id_z(); break;
+	default: *out = 0;
+	}
+}
+
+// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024}
+// CHECK-DAG: attributes #[[NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly }
+// CHECK-DAG: attributes #[[READ_EXEC_ATTRS]] = { convergent }
+// CHECK-DAG: ![[EXEC]] = !{!"exec"}
diff --git a/test/CodeGenOpenCL/builtins-generic-amdgcn.cl b/test/CodeGenOpenCL/builtins-generic-amdgcn.cl
new file mode 100644
index 000000000000..5a4756bacb85
--- /dev/null
+++ b/test/CodeGenOpenCL/builtins-generic-amdgcn.cl
@@ -0,0 +1,16 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
+
+// CHECK-LABEL: @test_builtin_clz(
+// CHECK: tail call i32 @llvm.ctlz.i32(i32 %a, i1 true)
+void test_builtin_clz(global int* out, int a)
+{
+  *out = __builtin_clz(a);
+}
+
+// CHECK-LABEL: @test_builtin_clzl(
+// CHECK: tail call i64 @llvm.ctlz.i64(i64 %a, i1 true)
+void test_builtin_clzl(global long* out, long a)
+{
+  *out = __builtin_clzl(a);
+}
diff --git a/test/CodeGenOpenCL/builtins-r600.cl b/test/CodeGenOpenCL/builtins-r600.cl
index 3e416b0323c2..027a54a6bce2 100644
--- a/test/CodeGenOpenCL/builtins-r600.cl
+++ b/test/CodeGenOpenCL/builtins-r600.cl
@@ -1,143 +1,55 @@
-// REQUIRES: r600-registered-target
-// RUN: %clang_cc1 -triple r600-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple r600-unknown-unknown -target-cpu cypress -S -emit-llvm -o - %s | FileCheck %s
 
-#pragma OPENCL EXTENSION cl_khr_fp64 : enable
-
-// CHECK-LABEL: @test_div_scale_f64
-// CHECK: call { double, i1 } @llvm.AMDGPU.div.scale.f64(double %a, double %b, i1 true)
-// CHECK-DAG: [[FLAG:%.+]] = extractvalue { double, i1 } %{{.+}}, 1
-// CHECK-DAG: [[VAL:%.+]] = extractvalue { double, i1 } %{{.+}}, 0
-// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i32
-// CHECK: store i32 [[FLAGEXT]]
-void test_div_scale_f64(global double* out, global int* flagout, double a, double b)
-{
-  bool flag;
-  *out = __builtin_amdgpu_div_scale(a, b, true, &flag);
-  *flagout = flag;
-}
-
-// CHECK-LABEL: @test_div_scale_f32
-// CHECK: call { float, i1 } @llvm.AMDGPU.div.scale.f32(float %a, float %b, i1 true)
-// CHECK-DAG: [[FLAG:%.+]] = extractvalue { float, i1 } %{{.+}}, 1
-// CHECK-DAG: [[VAL:%.+]] = extractvalue { float, i1 } %{{.+}}, 0
-// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i32
-// CHECK: store i32 [[FLAGEXT]]
-void test_div_scale_f32(global float* out, global int* flagout, float a, float b)
-{
-  bool flag;
-  *out = __builtin_amdgpu_div_scalef(a, b, true, &flag);
-  *flagout = flag;
-}
-
-// CHECK-LABEL: @test_div_fmas_f32
-// CHECK: call float @llvm.AMDGPU.div.fmas.f32
-void test_div_fmas_f32(global float* out, float a, float b, float c, int d)
-{
-  *out = __builtin_amdgpu_div_fmasf(a, b, c, d);
-}
-
-// CHECK-LABEL: @test_div_fmas_f64
-// CHECK: call double @llvm.AMDGPU.div.fmas.f64
-void test_div_fmas_f64(global double* out, double a, double b, double c, int d)
+// CHECK-LABEL: @test_recipsqrt_ieee_f32
+// CHECK: call float @llvm.r600.recipsqrt.ieee.f32
+void test_recipsqrt_ieee_f32(global float* out, float a)
 {
-  *out = __builtin_amdgpu_div_fmas(a, b, c, d);
+  *out = __builtin_r600_recipsqrt_ieeef(a);
 }
 
-// CHECK-LABEL: @test_div_fixup_f32
-// CHECK: call float @llvm.AMDGPU.div.fixup.f32
-void test_div_fixup_f32(global float* out, float a, float b, float c)
+#if cl_khr_fp64
+// XCHECK-LABEL: @test_recipsqrt_ieee_f64
+// XCHECK: call double @llvm.r600.recipsqrt.ieee.f64
+void test_recipsqrt_ieee_f64(global double* out, double a)
 {
-  *out = __builtin_amdgpu_div_fixupf(a, b, c);
+  *out = __builtin_r600_recipsqrt_ieee(a);
 }
+#endif
 
-// CHECK-LABEL: @test_div_fixup_f64
-// CHECK: call double @llvm.AMDGPU.div.fixup.f64
-void test_div_fixup_f64(global double* out, double a, double b, double c)
+// CHECK-LABEL: @test_implicitarg_ptr
+// CHECK: call i8 addrspace(7)* @llvm.r600.implicitarg.ptr()
+void test_implicitarg_ptr(__attribute__((address_space(7))) unsigned char ** out)
 {
-  *out = __builtin_amdgpu_div_fixup(a, b, c);
+  *out = __builtin_r600_implicitarg_ptr();
 }
 
-// CHECK-LABEL: @test_trig_preop_f32
-// CHECK: call float @llvm.AMDGPU.trig.preop.f32
-void test_trig_preop_f32(global float* out, float a, int b)
+// CHECK-LABEL: @test_get_group_id(
+// CHECK: tail call i32 @llvm.r600.read.tgid.x()
+// CHECK: tail call i32 @llvm.r600.read.tgid.y()
+// CHECK: tail call i32 @llvm.r600.read.tgid.z()
+void test_get_group_id(int d, global int *out)
 {
-  *out = __builtin_amdgpu_trig_preopf(a, b);
+	switch (d) {
+	case 0: *out = __builtin_r600_read_tgid_x(); break;
+	case 1: *out = __builtin_r600_read_tgid_y(); break;
+	case 2: *out = __builtin_r600_read_tgid_z(); break;
+	default: *out = 0;
+	}
 }
 
-// CHECK-LABEL: @test_trig_preop_f64
-// CHECK: call double @llvm.AMDGPU.trig.preop.f64
-void test_trig_preop_f64(global double* out, double a, int b)
+// CHECK-LABEL: @test_get_local_id(
+// CHECK: tail call i32 @llvm.r600.read.tidig.x(), !range [[WI_RANGE:![0-9]*]]
+// CHECK: tail call i32 @llvm.r600.read.tidig.y(), !range [[WI_RANGE]]
+// CHECK: tail call i32 @llvm.r600.read.tidig.z(), !range [[WI_RANGE]]
+void test_get_local_id(int d, global int *out)
 {
-  *out = __builtin_amdgpu_trig_preop(a, b);
+	switch (d) {
+	case 0: *out = __builtin_r600_read_tidig_x(); break;
+	case 1: *out = __builtin_r600_read_tidig_y(); break;
+	case 2: *out = __builtin_r600_read_tidig_z(); break;
+	default: *out = 0;
+	}
 }
 
-// CHECK-LABEL: @test_rcp_f32
-// CHECK: call float @llvm.AMDGPU.rcp.f32
-void test_rcp_f32(global float* out, float a)
-{
-  *out = __builtin_amdgpu_rcpf(a);
-}
-
-// CHECK-LABEL: @test_rcp_f64
-// CHECK: call double @llvm.AMDGPU.rcp.f64
-void test_rcp_f64(global double* out, double a)
-{
-  *out = __builtin_amdgpu_rcp(a);
-}
-
-// CHECK-LABEL: @test_rsq_f32
-// CHECK: call float @llvm.AMDGPU.rsq.f32
-void test_rsq_f32(global float* out, float a)
-{
-  *out = __builtin_amdgpu_rsqf(a);
-}
-
-// CHECK-LABEL: @test_rsq_f64
-// CHECK: call double @llvm.AMDGPU.rsq.f64
-void test_rsq_f64(global double* out, double a)
-{
-  *out = __builtin_amdgpu_rsq(a);
-}
-
-// CHECK-LABEL: @test_rsq_clamped_f32
-// CHECK: call float @llvm.AMDGPU.rsq.clamped.f32
-void test_rsq_clamped_f32(global float* out, float a)
-{
-  *out = __builtin_amdgpu_rsq_clampedf(a);
-}
-
-// CHECK-LABEL: @test_rsq_clamped_f64
-// CHECK: call double @llvm.AMDGPU.rsq.clamped.f64
-void test_rsq_clamped_f64(global double* out, double a)
-{
-  *out = __builtin_amdgpu_rsq_clamped(a);
-}
-
-// CHECK-LABEL: @test_ldexp_f32
-// CHECK: call float @llvm.AMDGPU.ldexp.f32
-void test_ldexp_f32(global float* out, float a, int b)
-{
-  *out = __builtin_amdgpu_ldexpf(a, b);
-}
-
-// CHECK-LABEL: @test_ldexp_f64
-// CHECK: call double @llvm.AMDGPU.ldexp.f64
-void test_ldexp_f64(global double* out, double a, int b)
-{
-  *out = __builtin_amdgpu_ldexp(a, b);
-}
-
-// CHECK-LABEL: @test_class_f32
-// CHECK: call i1 @llvm.AMDGPU.class.f32
-void test_class_f32(global float* out, float a, int b)
-{
-  *out = __builtin_amdgpu_classf(a, b);
-}
-
-// CHECK-LABEL: @test_class_f64
-// CHECK: call i1 @llvm.AMDGPU.class.f64
-void test_class_f64(global double* out, double a, int b)
-{
-  *out = __builtin_amdgpu_class(a, b);
-}
+// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024}
diff --git a/test/CodeGenOpenCL/cl20-device-side-enqueue.cl b/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
new file mode 100644
index 000000000000..08b24bdc5bec
--- /dev/null
+++ b/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
@@ -0,0 +1,110 @@
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - | FileCheck %s
+
+typedef void (^bl_t)(local void *);
+
+const bl_t block_G = (bl_t) ^ (local void *a) {};
+
+kernel void device_side_enqueue(global int *a, global int *b, int i) {
+  // CHECK: %default_queue = alloca %opencl.queue_t*
+  queue_t default_queue;
+  // CHECK: %flags = alloca i32
+  unsigned flags = 0;
+  // CHECK: %ndrange = alloca %opencl.ndrange_t*
+  ndrange_t ndrange;
+  // CHECK: %clk_event = alloca %opencl.clk_event_t*
+  clk_event_t clk_event;
+  // CHECK: %event_wait_list = alloca %opencl.clk_event_t*
+  clk_event_t event_wait_list;
+  // CHECK: %event_wait_list2 = alloca [1 x %opencl.clk_event_t*]
+  clk_event_t event_wait_list2[] = {clk_event};
+
+  // CHECK: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t*, %opencl.queue_t** %default_queue
+  // CHECK: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
+  // CHECK: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange
+  // CHECK: [[BL:%[0-9]+]] = bitcast <{ i8*, i32, i32, i8*, %struct.__block_descriptor*, i32{{.*}}, i32{{.*}}, i32{{.*}} }>* %block to void ()*
+  // CHECK: [[BL_I8:%[0-9]+]] = bitcast void ()* [[BL]] to i8*
+  // CHECK: call i32 @__enqueue_kernel_basic(%opencl.queue_t* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* [[BL_I8]])
+  enqueue_kernel(default_queue, flags, ndrange,
+                 ^(void) {
+                   a[i] = b[i];
+                 });
+
+  // CHECK: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t*, %opencl.queue_t** %default_queue
+  // CHECK: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
+  // CHECK: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange
+  // CHECK: [[BL:%[0-9]+]] = bitcast <{ i8*, i32, i32, i8*, %struct.__block_descriptor*, i32{{.*}}, i32{{.*}}, i32{{.*}} }>* %block3 to void ()*
+  // CHECK: [[BL_I8:%[0-9]+]] = bitcast void ()* [[BL]] to i8*
+  // CHECK: call i32 @__enqueue_kernel_basic_events(%opencl.queue_t* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i32 2, %opencl.clk_event_t** %event_wait_list, %opencl.clk_event_t** %clk_event, i8* [[BL_I8]])
+  enqueue_kernel(default_queue, flags, ndrange, 2, &event_wait_list, &clk_event,
+                 ^(void) {
+                   a[i] = b[i];
+                 });
+
+  // CHECK: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t*, %opencl.queue_t** %default_queue
+  // CHECK: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
+  // CHECK: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange
+  // CHECK: call i32 (%opencl.queue_t*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 256)
+  enqueue_kernel(default_queue, flags, ndrange,
+                 ^(local void *p) {
+                   return;
+                 },
+                 256);
+  char c;
+  // CHECK: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t*, %opencl.queue_t** %default_queue
+  // CHECK: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
+  // CHECK: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange
+  // CHECK: [[SIZE:%[0-9]+]] = zext i8 {{%[0-9]+}} to i32
+  // CHECK: call i32 (%opencl.queue_t*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 [[SIZE]])
+  enqueue_kernel(default_queue, flags, ndrange,
+                 ^(local void *p) {
+                   return;
+                 },
+                 c);
+
+  // CHECK: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t*, %opencl.queue_t** %default_queue
+  // CHECK: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
+  // CHECK: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange
+  // CHECK: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x %opencl.clk_event_t*], [1 x %opencl.clk_event_t*]* %event_wait_list2, i32 0, i32 0
+  // CHECK: call i32 (%opencl.queue_t*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t**, %opencl.clk_event_t**, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i32 2, %opencl.clk_event_t** [[AD]], %opencl.clk_event_t** %clk_event, i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 256)
+  enqueue_kernel(default_queue, flags, ndrange, 2, event_wait_list2, &clk_event,
+                 ^(local void *p) {
+                   return;
+                 },
+                 256);
+
+  // CHECK: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t*, %opencl.queue_t** %default_queue
+  // CHECK: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
+  // CHECK: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange
+  // CHECK: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x %opencl.clk_event_t*], [1 x %opencl.clk_event_t*]* %event_wait_list2, i32 0, i32 0
+  // CHECK: [[SIZE:%[0-9]+]] = zext i8 {{%[0-9]+}} to i32
+  // CHECK: call i32 (%opencl.queue_t*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t**, %opencl.clk_event_t**, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i32 2, %opencl.clk_event_t** [[AD]], %opencl.clk_event_t** %clk_event, i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 [[SIZE]])
+  enqueue_kernel(default_queue, flags, ndrange, 2, event_wait_list2, &clk_event,
+                 ^(local void *p) {
+                   return;
+                 },
+                 c);
+
+  void (^const block_A)(void) = ^{
+    return;
+  };
+  void (^const block_B)(local void *) = ^(local void *a) {
+    return;
+  };
+
+  // CHECK: [[BL:%[0-9]+]] = load void ()*, void ()** %block_A
+  // CHECK: [[BL_I8:%[0-9]+]] = bitcast void ()* [[BL]] to i8*
+  // CHECK: call i32 @__get_kernel_work_group_size_impl(i8* [[BL_I8]])
+  unsigned size = get_kernel_work_group_size(block_A);
+  // CHECK: [[BL:%[0-9]+]] = load void (i8 addrspace(2)*)*, void (i8 addrspace(2)*)** %block_B
+  // CHECK: [[BL_I8:%[0-9]+]] = bitcast void (i8 addrspace(2)*)* [[BL]] to i8*
+  // CHECK: call i32 @__get_kernel_work_group_size_impl(i8* [[BL_I8]])
+  size = get_kernel_work_group_size(block_B);
+  // CHECK: [[BL:%[0-9]+]] = load void ()*, void ()** %block_A
+  // CHECK: [[BL_I8:%[0-9]+]] = bitcast void ()* [[BL]] to i8*
+  // CHECK: call i32 @__get_kernel_preferred_work_group_multiple_impl(i8* [[BL_I8]])
+  size = get_kernel_preferred_work_group_size_multiple(block_A);
+  // CHECK: [[BL:%[0-9]+]] = load void (i8 addrspace(2)*)*, void (i8 addrspace(2)*)* addrspace(1)* @block_G
+  // CHECK: [[BL_I8:%[0-9]+]] = bitcast void (i8 addrspace(2)*)* [[BL]] to i8*
+  // CHECK: call i32 @__get_kernel_preferred_work_group_multiple_impl(i8* [[BL_I8]])
+  size = get_kernel_preferred_work_group_size_multiple(block_G);
+}
diff --git a/test/CodeGenOpenCL/constant-addr-space-globals.cl b/test/CodeGenOpenCL/constant-addr-space-globals.cl
index 92fb9790b5fb..f81a51458c42 100644
--- a/test/CodeGenOpenCL/constant-addr-space-globals.cl
+++ b/test/CodeGenOpenCL/constant-addr-space-globals.cl
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -cl-opt-disable -ffake-address-space-map -emit-llvm -o - | FileCheck %s
 
 // CHECK: @array = addrspace({{[0-9]+}}) constant
 __constant float array[2] = {0.0f, 1.0f};
diff --git a/test/CodeGenOpenCL/event_t.cl b/test/CodeGenOpenCL/event_t.cl
index a84d8bb610c0..aad441f35fd7 100644
--- a/test/CodeGenOpenCL/event_t.cl
+++ b/test/CodeGenOpenCL/event_t.cl
@@ -9,4 +9,6 @@ void kernel ker() {
 // CHECK: call {{.*}}void @foo(%opencl.event_t* %
   foo(0);
 // CHECK: call {{.*}}void @foo(%opencl.event_t* null)
+  foo((event_t)0);
+// CHECK: call {{.*}}void @foo(%opencl.event_t* null)
 }
diff --git a/test/CodeGenOpenCL/fpmath.cl b/test/CodeGenOpenCL/fpmath.cl
index ef4da845529c..88df3bf166ab 100644
--- a/test/CodeGenOpenCL/fpmath.cl
+++ b/test/CodeGenOpenCL/fpmath.cl
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
 
 typedef __attribute__(( ext_vector_type(4) )) float float4;
 
diff --git a/test/CodeGenOpenCL/half.cl b/test/CodeGenOpenCL/half.cl
index bd5ae7f64990..9acabf0a2a83 100644
--- a/test/CodeGenOpenCL/half.cl
+++ b/test/CodeGenOpenCL/half.cl
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
diff --git a/test/CodeGenOpenCL/images.cl b/test/CodeGenOpenCL/images.cl
new file mode 100644
index 000000000000..eb054eceb5df
--- /dev/null
+++ b/test/CodeGenOpenCL/images.cl
@@ -0,0 +1,11 @@
+// RUN: %clang_cc1 %s -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -o - | FileCheck %s
+
+__attribute__((overloadable)) void read_image(read_only image1d_t img_ro);
+__attribute__((overloadable)) void read_image(write_only image1d_t img_wo);
+
+kernel void test_read_image(read_only image1d_t img_ro, write_only image1d_t img_wo) {
+  // CHECK: call void @_Z10read_image14ocl_image1d_ro(%opencl.image1d_ro_t* %{{[0-9]+}})
+  read_image(img_ro);
+  // CHECK: call void @_Z10read_image14ocl_image1d_wo(%opencl.image1d_wo_t* %{{[0-9]+}})
+  read_image(img_wo);
+}
diff --git a/test/CodeGenOpenCL/kernel-arg-info.cl b/test/CodeGenOpenCL/kernel-arg-info.cl
index 4bc191e1d75f..5a5c8f9d6c5f 100644
--- a/test/CodeGenOpenCL/kernel-arg-info.cl
+++ b/test/CodeGenOpenCL/kernel-arg-info.cl
@@ -1,55 +1,88 @@
-// RUN: %clang_cc1 %s -cl-kernel-arg-info -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s -check-prefix ARGINFO
-// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s -check-prefix NO-ARGINFO
+// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown -cl-kernel-arg-info | FileCheck %s -check-prefix ARGINFO
 
 kernel void foo(__global int * restrict X, const int Y, 
                 volatile int anotherArg, __constant float * restrict Z) {
   *X = Y + anotherArg;
 }
-
-// CHECK:  !{!"kernel_arg_addr_space", i32 1, i32 0, i32 0, i32 2}
-// CHECK:  !{!"kernel_arg_access_qual", !"none", !"none", !"none", !"none"}
-// CHECK:  !{!"kernel_arg_type", !"int*", !"int", !"int", !"float*"}
-// CHECK:  !{!"kernel_arg_base_type", !"int*", !"int", !"int", !"float*"}
-// CHECK:  !{!"kernel_arg_type_qual", !"restrict", !"const", !"volatile", !"restrict const"}
-// ARGINFO: !{!"kernel_arg_name", !"X", !"Y", !"anotherArg", !"Z"}
-// NO-ARGINFO-NOT: !{!"kernel_arg_name", !"X", !"Y", !"anotherArg", !"Z"}
+// CHECK: define spir_kernel void @foo{{[^!]+}}
+// CHECK: !kernel_arg_addr_space ![[MD11:[0-9]+]]
+// CHECK: !kernel_arg_access_qual ![[MD12:[0-9]+]]
+// CHECK: !kernel_arg_type ![[MD13:[0-9]+]]
+// CHECK: !kernel_arg_base_type ![[MD13]]
+// CHECK: !kernel_arg_type_qual ![[MD14:[0-9]+]]
+// CHECK-NOT: !kernel_arg_name
+// ARGINFO: !kernel_arg_name ![[MD15:[0-9]+]]
 
 kernel void foo2(read_only image1d_t img1, image2d_t img2, write_only image2d_array_t img3) {
 }
-// CHECK:  !{!"kernel_arg_addr_space", i32 1, i32 1, i32 1}
-// CHECK:  !{!"kernel_arg_access_qual", !"read_only", !"read_only", !"write_only"}
-// CHECK:  !{!"kernel_arg_type", !"image1d_t", !"image2d_t", !"image2d_array_t"}
-// CHECK:  !{!"kernel_arg_base_type", !"image1d_t", !"image2d_t", !"image2d_array_t"}
-// CHECK:  !{!"kernel_arg_type_qual", !"", !"", !""}
-// ARGINFO: !{!"kernel_arg_name", !"img1", !"img2", !"img3"}
-// NO-ARGINFO-NOT: !{!"kernel_arg_name", !"img1", !"img2", !"img3"}
+// CHECK: define spir_kernel void @foo2{{[^!]+}}
+// CHECK: !kernel_arg_addr_space ![[MD21:[0-9]+]]
+// CHECK: !kernel_arg_access_qual ![[MD22:[0-9]+]]
+// CHECK: !kernel_arg_type ![[MD23:[0-9]+]]
+// CHECK: !kernel_arg_base_type ![[MD23]]
+// CHECK: !kernel_arg_type_qual ![[MD24:[0-9]+]]
+// CHECK-NOT: !kernel_arg_name
+// ARGINFO: !kernel_arg_name ![[MD25:[0-9]+]]
 
 kernel void foo3(__global half * X) {
 }
-// CHECK:  !{!"kernel_arg_addr_space", i32 1}
-// CHECK:  !{!"kernel_arg_access_qual", !"none"}
-// CHECK:  !{!"kernel_arg_type", !"half*"}
-// CHECK:  !{!"kernel_arg_base_type", !"half*"}
-// CHECK:  !{!"kernel_arg_type_qual", !""}
-// ARGINFO: !{!"kernel_arg_name", !"X"}
-// NO-ARGINFO-NOT: !{!"kernel_arg_name", !"X"}
+// CHECK: define spir_kernel void @foo3{{[^!]+}}
+// CHECK: !kernel_arg_addr_space ![[MD31:[0-9]+]]
+// CHECK: !kernel_arg_access_qual ![[MD32:[0-9]+]]
+// CHECK: !kernel_arg_type ![[MD33:[0-9]+]]
+// CHECK: !kernel_arg_base_type ![[MD33]]
+// CHECK: !kernel_arg_type_qual ![[MD34:[0-9]+]]
+// CHECK-NOT: !kernel_arg_name
+// ARGINFO: !kernel_arg_name ![[MD35:[0-9]+]]
 
 typedef unsigned int myunsignedint;
 kernel void foo4(__global unsigned int * X, __global myunsignedint * Y) {
 }
-// CHECK:  !{!"kernel_arg_addr_space", i32 1, i32 1}
-// CHECK:  !{!"kernel_arg_access_qual", !"none", !"none"}
-// CHECK:  !{!"kernel_arg_type", !"uint*", !"myunsignedint*"}
-// CHECK:  !{!"kernel_arg_base_type", !"uint*", !"uint*"}
-// CHECK:  !{!"kernel_arg_type_qual", !"", !""}
-// ARGINFO: !{!"kernel_arg_name", !"X", !"Y"}
-// NO-ARGINFO-NOT: !{!"kernel_arg_name", !"X", !"Y"}
+// CHECK: define spir_kernel void @foo4{{[^!]+}}
+// CHECK: !kernel_arg_addr_space ![[MD41:[0-9]+]]
+// CHECK: !kernel_arg_access_qual ![[MD42:[0-9]+]]
+// CHECK: !kernel_arg_type ![[MD43:[0-9]+]]
+// CHECK: !kernel_arg_base_type ![[MD44:[0-9]+]]
+// CHECK: !kernel_arg_type_qual ![[MD45:[0-9]+]]
+// CHECK-NOT: !kernel_arg_name
+// ARGINFO: !kernel_arg_name ![[MD46:[0-9]+]]
 
 typedef image1d_t myImage;
-kernel void foo5(read_only myImage img1, write_only image1d_t img2) {
+kernel void foo5(myImage img1, write_only image1d_t img2) {
 }
-// CHECK:  !{!"kernel_arg_access_qual", !"read_only", !"write_only"}
-// CHECK:  !{!"kernel_arg_type", !"myImage", !"image1d_t"}
-// CHECK:  !{!"kernel_arg_base_type", !"image1d_t", !"image1d_t"}
-// ARGINFO: !{!"kernel_arg_name", !"img1", !"img2"}
-// NO-ARGINFO-NOT: !{!"kernel_arg_name", !"img1", !"img2"}
+// CHECK: define spir_kernel void @foo5{{[^!]+}}
+// CHECK: !kernel_arg_addr_space ![[MD41:[0-9]+]]
+// CHECK: !kernel_arg_access_qual ![[MD51:[0-9]+]]
+// CHECK: !kernel_arg_type ![[MD52:[0-9]+]]
+// CHECK: !kernel_arg_base_type ![[MD53:[0-9]+]]
+// CHECK: !kernel_arg_type_qual ![[MD45]]
+// CHECK-NOT: !kernel_arg_name
+// ARGINFO: !kernel_arg_name ![[MD54:[0-9]+]]
+
+// CHECK: ![[MD11]] = !{i32 1, i32 0, i32 0, i32 2}
+// CHECK: ![[MD12]] = !{!"none", !"none", !"none", !"none"}
+// CHECK: ![[MD13]] = !{!"int*", !"int", !"int", !"float*"}
+// CHECK: ![[MD14]] = !{!"restrict", !"const", !"volatile", !"restrict const"}
+// ARGINFO: ![[MD15]] = !{!"X", !"Y", !"anotherArg", !"Z"}
+// CHECK: ![[MD21]] = !{i32 1, i32 1, i32 1}
+// CHECK: ![[MD22]] = !{!"read_only", !"read_only", !"write_only"}
+// CHECK: ![[MD23]] = !{!"__read_only image1d_t", !"__read_only image2d_t", !"__write_only image2d_array_t"}
+// CHECK: ![[MD24]] = !{!"", !"", !""}
+// ARGINFO: ![[MD25]] = !{!"img1", !"img2", !"img3"}
+// CHECK: ![[MD31]] = !{i32 1}
+// CHECK: ![[MD32]] = !{!"none"}
+// CHECK: ![[MD33]] = !{!"half*"}
+// CHECK: ![[MD34]] = !{!""}
+// ARGINFO: ![[MD35]] = !{!"X"}
+// CHECK: ![[MD41]] = !{i32 1, i32 1}
+// CHECK: ![[MD42]] = !{!"none", !"none"}
+// CHECK: ![[MD43]] = !{!"uint*", !"myunsignedint*"}
+// CHECK: ![[MD44]] = !{!"uint*", !"uint*"}
+// CHECK: ![[MD45]] = !{!"", !""}
+// ARGINFO: ![[MD46]] = !{!"X", !"Y"}
+// CHECK: ![[MD51]] = !{!"read_only", !"write_only"}
+// CHECK: ![[MD52]] = !{!"myImage", !"__write_only image1d_t"}
+// CHECK: ![[MD53]] = !{!"__read_only image1d_t", !"__write_only image1d_t"}
+// ARGINFO: ![[MD54]] = !{!"img1", !"img2"}
+
diff --git a/test/CodeGenOpenCL/kernel-attributes.cl b/test/CodeGenOpenCL/kernel-attributes.cl
index 8f22d611b820..4a116dd08551 100644
--- a/test/CodeGenOpenCL/kernel-attributes.cl
+++ b/test/CodeGenOpenCL/kernel-attributes.cl
@@ -3,14 +3,12 @@
 typedef unsigned int uint4 __attribute__((ext_vector_type(4)));
 
 kernel  __attribute__((vec_type_hint(int))) __attribute__((reqd_work_group_size(1,2,4))) void kernel1(int a) {}
+// CHECK: define void @kernel1(i32 {{[^%]*}}%a) {{[^{]+}} !vec_type_hint ![[MD1:[0-9]+]] !reqd_work_group_size ![[MD2:[0-9]+]]
 
 kernel __attribute__((vec_type_hint(uint4))) __attribute__((work_group_size_hint(8,16,32))) void kernel2(int a) {}
+// CHECK: define void @kernel2(i32 {{[^%]*}}%a) {{[^{]+}} !vec_type_hint ![[MD3:[0-9]+]] !work_group_size_hint ![[MD4:[0-9]+]]
 
-// CHECK: opencl.kernels = !{[[MDNODE0:![0-9]+]], [[MDNODE3:![0-9]+]]}
-
-// CHECK: [[MDNODE0]] = !{void (i32)* @kernel1, {{.*}} [[MDNODE1:![0-9]+]], [[MDNODE2:![0-9]+]]}
-// CHECK: [[MDNODE1]] = !{!"vec_type_hint", i32 undef, i32 1}
-// CHECK: [[MDNODE2]] = !{!"reqd_work_group_size", i32 1, i32 2, i32 4}
-// CHECK: [[MDNODE3]] = !{void (i32)* @kernel2, {{.*}} [[MDNODE4:![0-9]+]], [[MDNODE5:![0-9]+]]}
-// CHECK: [[MDNODE4]] = !{!"vec_type_hint", <4 x i32> undef, i32 0}
-// CHECK: [[MDNODE5]] = !{!"work_group_size_hint", i32 8, i32 16, i32 32}
+// CHECK: [[MD1]] = !{i32 undef, i32 1}
+// CHECK: [[MD2]] = !{i32 1, i32 2, i32 4}
+// CHECK: [[MD3]] = !{<4 x i32> undef, i32 0}
+// CHECK: [[MD4]] = !{i32 8, i32 16, i32 32}
diff --git a/test/CodeGenOpenCL/kernel-metadata.cl b/test/CodeGenOpenCL/kernel-metadata.cl
index ef3758fccaa0..4165f1fa0ce5 100644
--- a/test/CodeGenOpenCL/kernel-metadata.cl
+++ b/test/CodeGenOpenCL/kernel-metadata.cl
@@ -6,10 +6,5 @@ void normal_function() {
 __kernel void kernel_function() {
 }
 
-// CHECK: !opencl.kernels = !{!0}
-// CHECK: !0 = !{void ()* @kernel_function, !1, !2, !3, !4, !5}
-// CHECK: !1 = !{!"kernel_arg_addr_space"}
-// CHECK: !2 = !{!"kernel_arg_access_qual"}
-// CHECK: !3 = !{!"kernel_arg_type"}
-// CHECK: !4 = !{!"kernel_arg_base_type"}
-// CHECK: !5 = !{!"kernel_arg_type_qual"}
+// CHECK: define void @kernel_function() {{[^{]+}} !kernel_arg_addr_space ![[MD:[0-9]+]] !kernel_arg_access_qual ![[MD]] !kernel_arg_type ![[MD]] !kernel_arg_base_type ![[MD]] !kernel_arg_type_qual ![[MD]] {
+// CHECK: ![[MD]] = !{}
diff --git a/test/CodeGenOpenCL/no-signed-zeros.cl b/test/CodeGenOpenCL/no-signed-zeros.cl
new file mode 100644
index 000000000000..14f6411b3561
--- /dev/null
+++ b/test/CodeGenOpenCL/no-signed-zeros.cl
@@ -0,0 +1,10 @@
+// RUN: %clang_cc1 %s -emit-llvm -o - | FileCheck %s -check-prefix=NORMAL

+// RUN: %clang_cc1 %s -emit-llvm -cl-no-signed-zeros -o - | FileCheck %s -check-prefix=NO-SIGNED-ZEROS

+

+float signedzeros(float a) {

+  return a;

+}

+

+// CHECK: attributes

+// NORMAL: "no-signed-zeros-fp-math"="false"

+// NO-SIGNED-ZEROS: "no-signed-zeros-fp-math"="true"

diff --git a/test/CodeGenOpenCL/opencl_types.cl b/test/CodeGenOpenCL/opencl_types.cl
index 5f4ebb8a283e..5691b27fbe52 100644
--- a/test/CodeGenOpenCL/opencl_types.cl
+++ b/test/CodeGenOpenCL/opencl_types.cl
@@ -4,37 +4,37 @@ constant sampler_t glb_smp = 7;
 // CHECK: constant i32 7
 
 void fnc1(image1d_t img) {}
-// CHECK: @fnc1(%opencl.image1d_t*
+// CHECK: @fnc1(%opencl.image1d_ro_t*
 
 void fnc1arr(image1d_array_t img) {}
-// CHECK: @fnc1arr(%opencl.image1d_array_t*
+// CHECK: @fnc1arr(%opencl.image1d_array_ro_t*
 
 void fnc1buff(image1d_buffer_t img) {}
-// CHECK: @fnc1buff(%opencl.image1d_buffer_t*
+// CHECK: @fnc1buff(%opencl.image1d_buffer_ro_t*
 
 void fnc2(image2d_t img) {}
-// CHECK: @fnc2(%opencl.image2d_t*
+// CHECK: @fnc2(%opencl.image2d_ro_t*
 
 void fnc2arr(image2d_array_t img) {}
-// CHECK: @fnc2arr(%opencl.image2d_array_t*
+// CHECK: @fnc2arr(%opencl.image2d_array_ro_t*
 
 void fnc3(image3d_t img) {}
-// CHECK: @fnc3(%opencl.image3d_t*
+// CHECK: @fnc3(%opencl.image3d_ro_t*
 
 void fnc4smp(sampler_t s) {}
 // CHECK-LABEL: define {{.*}}void @fnc4smp(i32
 
 kernel void foo(image1d_t img) {
-	sampler_t smp = 5;
-// CHECK: alloca i32
-	event_t evt;
-// CHECK: alloca %opencl.event_t*
-// CHECK: store i32 5,
+  sampler_t smp = 5;
+  // CHECK: alloca i32
+  event_t evt;
+  // CHECK: alloca %opencl.event_t*
+  // CHECK: store i32 5,
   fnc4smp(smp);
-// CHECK: call {{.*}}void @fnc4smp(i32
+  // CHECK: call {{.*}}void @fnc4smp(i32
   fnc4smp(glb_smp);
-// CHECK: call {{.*}}void @fnc4smp(i32
+  // CHECK: call {{.*}}void @fnc4smp(i32
 }
 
-void __attribute__((overloadable)) bad1(image1d_t *b, image2d_t *c, image2d_t *d) {}
-// CHECK-LABEL: @{{_Z4bad1P11ocl_image1dP11ocl_image2dS2_|"\\01\?bad1@@\$\$J0YAXPE?APAUocl_image1d@@PE?APAUocl_image2d@@1@Z"}}
+void __attribute__((overloadable)) bad1(image1d_t b, image2d_t c, image2d_t d) {}
+// CHECK-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}}
diff --git a/test/CodeGenOpenCL/pipe_builtin.cl b/test/CodeGenOpenCL/pipe_builtin.cl
new file mode 100644
index 000000000000..db6893ebf55d
--- /dev/null
+++ b/test/CodeGenOpenCL/pipe_builtin.cl
@@ -0,0 +1,61 @@
+// RUN: %clang_cc1 -emit-llvm -O0 -cl-std=CL2.0 -o - %s | FileCheck %s
+
+// CHECK: %opencl.pipe_t = type opaque
+// CHECK: %opencl.reserve_id_t = type opaque
+
+void test1(read_only pipe int p, global int *ptr) {
+  // CHECK: call i32 @__read_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}})
+  read_pipe(p, ptr);
+  // CHECK: call %opencl.reserve_id_t* @__reserve_read_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}})
+  reserve_id_t rid = reserve_read_pipe(p, 2);
+  // CHECK: call i32 @__read_pipe_4(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 {{.*}}, i8* %{{.*}})
+  read_pipe(p, rid, 2, ptr);
+  // CHECK: call void @__commit_read_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}})
+  commit_read_pipe(p, rid);
+}
+
+void test2(write_only pipe int p, global int *ptr) {
+  // CHECK: call i32 @__write_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}})
+  write_pipe(p, ptr);
+  // CHECK: call %opencl.reserve_id_t* @__reserve_write_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}})
+  reserve_id_t rid = reserve_write_pipe(p, 2);
+  // CHECK: call i32 @__write_pipe_4(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}}, i32 {{.*}}, i8* %{{.*}})
+  write_pipe(p, rid, 2, ptr);
+  // CHECK: call void @__commit_write_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}})
+  commit_write_pipe(p, rid);
+}
+
+void test3(read_only pipe int p, global int *ptr) {
+  // CHECK: call %opencl.reserve_id_t* @__work_group_reserve_read_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}})
+  reserve_id_t rid = work_group_reserve_read_pipe(p, 2);
+  // CHECK: call void @__work_group_commit_read_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}})
+  work_group_commit_read_pipe(p, rid);
+}
+
+void test4(write_only pipe int p, global int *ptr) {
+  // CHECK: call %opencl.reserve_id_t* @__work_group_reserve_write_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}})
+  reserve_id_t rid = work_group_reserve_write_pipe(p, 2);
+  // CHECK: call void @__work_group_commit_write_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}})
+  work_group_commit_write_pipe(p, rid);
+}
+
+void test5(read_only pipe int p, global int *ptr) {
+  // CHECK: call %opencl.reserve_id_t* @__sub_group_reserve_read_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}})
+  reserve_id_t rid = sub_group_reserve_read_pipe(p, 2);
+  // CHECK: call void @__sub_group_commit_read_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}})
+  sub_group_commit_read_pipe(p, rid);
+}
+
+void test6(write_only pipe int p, global int *ptr) {
+  // CHECK: call %opencl.reserve_id_t* @__sub_group_reserve_write_pipe(%opencl.pipe_t* %{{.*}}, i32 {{.*}})
+  reserve_id_t rid = sub_group_reserve_write_pipe(p, 2);
+  // CHECK: call void @__sub_group_commit_write_pipe(%opencl.pipe_t* %{{.*}}, %opencl.reserve_id_t* %{{.*}})
+  sub_group_commit_write_pipe(p, rid);
+}
+
+void test7(write_only pipe int p, global int *ptr) {
+  // CHECK: call i32 @__get_pipe_num_packets(%opencl.pipe_t* %{{.*}})
+  *ptr = get_pipe_num_packets(p);
+  // CHECK: call i32 @__get_pipe_max_packets(%opencl.pipe_t* %{{.*}})
+  *ptr = get_pipe_max_packets(p);
+}
diff --git a/test/CodeGenOpenCL/pipe_types.cl b/test/CodeGenOpenCL/pipe_types.cl
index 9d8d5274a3ba..b9c411b5fcb6 100644
--- a/test/CodeGenOpenCL/pipe_types.cl
+++ b/test/CodeGenOpenCL/pipe_types.cl
@@ -25,3 +25,23 @@ void test4(read_only pipe uchar3 p) {
 void test5(read_only pipe int4 p) {
 // CHECK: define void @test5(%opencl.pipe_t* %p)
 }
+
+typedef read_only pipe int MyPipe;
+kernel void test6(MyPipe p) {
+// CHECK: define void @test6(%opencl.pipe_t* %p)
+}
+
+struct Person {
+  const char *Name;
+  bool isFemale;
+  int ID;
+};
+
+void test_reserved_read_pipe(global struct Person *SDst,
+                             read_only pipe struct Person SPipe) {
+// CHECK: define void @test_reserved_read_pipe
+  read_pipe (SPipe, SDst);
+  // CHECK: call i32 @__read_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}})
+  read_pipe (SPipe, SDst);
+  // CHECK: call i32 @__read_pipe_2(%opencl.pipe_t* %{{.*}}, i8* %{{.*}})
+}
diff --git a/test/CodeGenOpenCL/shifts.cl b/test/CodeGenOpenCL/shifts.cl
index ab64051a01ce..14cd7aff65be 100644
--- a/test/CodeGenOpenCL/shifts.cl
+++ b/test/CodeGenOpenCL/shifts.cl
@@ -5,7 +5,7 @@
 // bits before evaluating. Test this both for variables and constants
 // evaluated in the front-end.
 
-// OPT: @gtest1 = constant i64 2147483648
+// OPT: @gtest1 = local_unnamed_addr constant i64 2147483648
 __constant const unsigned long gtest1 = 1UL << 31;
 
 // NOOPT: @negativeShift32
diff --git a/test/CodeGenOpenCL/spir_version.cl b/test/CodeGenOpenCL/spir_version.cl
new file mode 100644
index 000000000000..215b4d7a4023
--- /dev/null
+++ b/test/CodeGenOpenCL/spir_version.cl
@@ -0,0 +1,18 @@
+// RUN: %clang_cc1 %s -triple "spir-unknown-unknown" -emit-llvm -o - | FileCheck %s --check-prefix=CL10
+// RUN: %clang_cc1 %s -triple "spir-unknown-unknown" -emit-llvm -o - -cl-std=CL1.2 | FileCheck %s --check-prefix=CL12
+// RUN: %clang_cc1 %s -triple "spir-unknown-unknown" -emit-llvm -o - -cl-std=CL2.0 | FileCheck %s --check-prefix=CL20
+// RUN: %clang_cc1 %s -triple "spir64-unknown-unknown" -emit-llvm -o - | FileCheck %s --check-prefix=CL10
+// RUN: %clang_cc1 %s -triple "spir64-unknown-unknown" -emit-llvm -o - -cl-std=CL1.2 | FileCheck %s --check-prefix=CL12
+// RUN: %clang_cc1 %s -triple "spir64-unknown-unknown" -emit-llvm -o - -cl-std=CL2.0 | FileCheck %s --check-prefix=CL20
+kernel void foo() {}
+// CL10: !opencl.spir.version = !{[[SPIR:![0-9]+]]}
+// CL10: !opencl.ocl.version = !{[[OCL:![0-9]+]]}
+// CL10: [[SPIR]] = !{i32 2, i32 0}
+// CL10: [[OCL]] = !{i32 1, i32 0}
+// CL12: !opencl.spir.version = !{[[SPIR:![0-9]+]]}
+// CL12: !opencl.ocl.version = !{[[OCL:![0-9]+]]}
+// CL12: [[SPIR]] = !{i32 2, i32 0}
+// CL12: [[OCL]] = !{i32 1, i32 2}
+// CL20: !opencl.spir.version = !{[[SPIR:![0-9]+]]}
+// CL20: !opencl.ocl.version = !{[[SPIR:![0-9]+]]}
+// CL20: [[SPIR]] = !{i32 2, i32 0}
diff --git a/test/CodeGenOpenCL/str_literals.cl b/test/CodeGenOpenCL/str_literals.cl
index 092b6372a414..1c0acd155dac 100644
--- a/test/CodeGenOpenCL/str_literals.cl
+++ b/test/CodeGenOpenCL/str_literals.cl
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -emit-llvm -o - -ffake-address-space-map | FileCheck %s
+// RUN: %clang_cc1 %s -cl-opt-disable -emit-llvm -o - -ffake-address-space-map | FileCheck %s
 
 __constant char * __constant x = "hello world";
 __constant char * __constant y = "hello world";
diff --git a/test/CodeGenOpenCL/to_addr_builtin.cl b/test/CodeGenOpenCL/to_addr_builtin.cl
new file mode 100644
index 000000000000..67475d565e3a
--- /dev/null
+++ b/test/CodeGenOpenCL/to_addr_builtin.cl
@@ -0,0 +1,89 @@
+// RUN: %clang_cc1 -triple spir-unknown-unknown -emit-llvm -O0 -cl-std=CL2.0 -o - %s | FileCheck %s
+
+// CHECK: %[[A:.*]] = type { float, float, float }
+typedef struct {
+  float x,y,z;
+} A;
+typedef private A *PA;
+typedef global A *GA;
+
+void test(void) {
+  global int *glob;
+  local int *loc;
+  private int *priv;
+  generic int *gen;
+
+  //CHECK: %[[ARG:.*]] = addrspacecast i32 addrspace(1)* %{{.*}} to i8 addrspace(4)*
+  //CHECK: %[[RET:.*]] = call i8 addrspace(1)* @to_global(i8 addrspace(4)* %[[ARG]])
+  //CHECK: %{{.*}} = bitcast i8 addrspace(1)* %[[RET]] to i32 addrspace(1)*
+  glob = to_global(glob);
+  
+  //CHECK: %[[ARG:.*]] = addrspacecast i32 addrspace(3)* %{{.*}} to i8 addrspace(4)*
+  //CHECK: %[[RET:.*]] = call i8 addrspace(1)* @to_global(i8 addrspace(4)* %[[ARG]])
+  //CHECK: %{{.*}} = bitcast i8 addrspace(1)* %[[RET]] to i32 addrspace(1)*
+  glob = to_global(loc);
+ 
+  //CHECK: %[[ARG:.*]] = addrspacecast i32* %{{.*}} to i8 addrspace(4)*
+  //CHECK: %[[RET:.*]] = call i8 addrspace(1)* @to_global(i8 addrspace(4)* %[[ARG]])
+  //CHECK: %{{.*}} = bitcast i8 addrspace(1)* %[[RET]] to i32 addrspace(1)*
+  glob = to_global(priv);
+ 
+  //CHECK: %[[ARG:.*]] = bitcast i32 addrspace(4)* %{{.*}} to i8 addrspace(4)*
+  //CHECK: %[[RET:.*]] = call i8 addrspace(1)* @to_global(i8 addrspace(4)* %[[ARG]])
+  //CHECK: %{{.*}} = bitcast i8 addrspace(1)* %[[RET]] to i32 addrspace(1)*
+  glob = to_global(gen);
+  
+  //CHECK: %[[ARG:.*]] = addrspacecast i32 addrspace(1)* %{{.*}} to i8 addrspace(4)*
+  //CHECK: %[[RET:.*]] = call i8 addrspace(3)* @to_local(i8 addrspace(4)* %[[ARG]])
+  //CHECK: %{{.*}} = bitcast i8 addrspace(3)* %[[RET]] to i32 addrspace(3)*
+  loc = to_local(glob);
+
+  //CHECK: %[[ARG:.*]] = addrspacecast i32 addrspace(3)* %{{.*}} to i8 addrspace(4)*
+  //CHECK: %[[RET:.*]] = call i8 addrspace(3)* @to_local(i8 addrspace(4)* %[[ARG]])
+  //CHECK: %{{.*}} = bitcast i8 addrspace(3)* %[[RET]] to i32 addrspace(3)*
+  loc = to_local(loc);
+
+  //CHECK: %[[ARG:.*]] = addrspacecast i32* %{{.*}} to i8 addrspace(4)*
+  //CHECK: %[[RET:.*]] = call i8 addrspace(3)* @to_local(i8 addrspace(4)* %[[ARG]])
+  //CHECK: %{{.*}} = bitcast i8 addrspace(3)* %[[RET]] to i32 addrspace(3)*
+  loc = to_local(priv);
+
+  //CHECK: %[[ARG:.*]] = bitcast i32 addrspace(4)* %{{.*}} to i8 addrspace(4)*
+  //CHECK: %[[RET:.*]] = call i8 addrspace(3)* @to_local(i8 addrspace(4)* %[[ARG]])
+  //CHECK: %{{.*}} = bitcast i8 addrspace(3)* %[[RET]] to i32 addrspace(3)*
+  loc = to_local(gen);
+
+  //CHECK: %[[ARG:.*]] = addrspacecast i32 addrspace(1)* %{{.*}} to i8 addrspace(4)*
+  //CHECK: %[[RET:.*]] = call i8* @to_private(i8 addrspace(4)* %[[ARG]])
+  //CHECK: %{{.*}} = bitcast i8* %[[RET]] to i32*
+  priv = to_private(glob);
+
+  //CHECK: %[[ARG:.*]] = addrspacecast i32 addrspace(3)* %{{.*}} to i8 addrspace(4)*
+  //CHECK: %[[RET:.*]] = call i8* @to_private(i8 addrspace(4)* %[[ARG]])
+  //CHECK: %{{.*}} = bitcast i8* %[[RET]] to i32*
+  priv = to_private(loc);
+
+  //CHECK: %[[ARG:.*]] = addrspacecast i32* %{{.*}} to i8 addrspace(4)*
+  //CHECK: %[[RET:.*]] = call i8* @to_private(i8 addrspace(4)* %[[ARG]])
+  //CHECK: %{{.*}} = bitcast i8* %[[RET]] to i32*
+  priv = to_private(priv);
+
+  //CHECK: %[[ARG:.*]] = bitcast i32 addrspace(4)* %{{.*}} to i8 addrspace(4)*
+  //CHECK: %[[RET:.*]] = call i8* @to_private(i8 addrspace(4)* %[[ARG]])
+  //CHECK: %{{.*}} = bitcast i8* %[[RET]] to i32*
+  priv = to_private(gen);
+
+  //CHECK: %[[ARG:.*]] = addrspacecast %[[A]]* %{{.*}} to i8 addrspace(4)*
+  //CHECK: %[[RET:.*]] = call i8 addrspace(1)* @to_global(i8 addrspace(4)* %[[ARG]])
+  //CHECK: %{{.*}} = bitcast i8 addrspace(1)* %[[RET]] to %[[A]] addrspace(1)*
+  PA pA;
+  GA gA = to_global(pA);
+
+  //CHECK-NOT: addrspacecast
+  //CHECK-NOT: bitcast
+  //CHECK: call i8 addrspace(1)* @to_global(i8 addrspace(4)* %{{.*}})
+  //CHECK-NOT: addrspacecast
+  //CHECK-NOT: bitcast
+  generic void *gen_v;
+  global void *glob_v = to_global(gen_v);
+}
diff --git a/test/CodeGenOpenCL/unroll-hint.cl b/test/CodeGenOpenCL/unroll-hint.cl
new file mode 100644
index 000000000000..a86762e02b54
--- /dev/null
+++ b/test/CodeGenOpenCL/unroll-hint.cl
@@ -0,0 +1,96 @@
+// RUN: %clang_cc1 -emit-llvm -O0 -cl-std=CL2.0 -o - %s | FileCheck %s
+
+/*** for ***/
+void for_count()
+{
+// CHECK-LABEL: for_count
+    __attribute__((opencl_unroll_hint(8)))
+    for( int i = 0; i < 1000; ++i);
+// CHECK: br label %{{.*}}, !llvm.loop ![[FOR_COUNT:.*]]
+}
+
+void for_disable()
+{
+// CHECK-LABEL: for_disable
+    __attribute__((opencl_unroll_hint(1)))
+    for( int i = 0; i < 1000; ++i);
+// CHECK: br label %{{.*}}, !llvm.loop ![[FOR_DISABLE:.*]]
+}
+
+void for_full()
+{
+// CHECK-LABEL: for_full
+    __attribute__((opencl_unroll_hint))
+    for( int i = 0; i < 1000; ++i);
+// CHECK: br label %{{.*}}, !llvm.loop ![[FOR_FULL:.*]]
+}
+
+/*** while ***/
+void while_count()
+{
+// CHECK-LABEL: while_count
+    int i = 1000;
+    __attribute__((opencl_unroll_hint(8)))
+    while(i-->0);
+// CHECK: br label %{{.*}}, !llvm.loop ![[WHILE_COUNT:.*]]
+}
+
+void while_disable()
+{
+// CHECK-LABEL: while_disable
+    int i = 1000;
+    __attribute__((opencl_unroll_hint(1)))
+    while(i-->0);
+// CHECK: br label %{{.*}}, !llvm.loop ![[WHILE_DISABLE:.*]]
+}
+
+void while_full()
+{
+// CHECK-LABEL: while_full
+    int i = 1000;
+    __attribute__((opencl_unroll_hint))
+    while(i-->0);
+// CHECK: br label %{{.*}}, !llvm.loop ![[WHILE_FULL:.*]]
+}
+
+/*** do ***/
+void do_count()
+{
+// CHECK-LABEL: do_count
+    int i = 1000;
+    __attribute__((opencl_unroll_hint(8)))
+    do {} while(i--> 0);
+// CHECK: br i1 %{{.*}}, label %{{.*}}, label %{{.*}}, !llvm.loop ![[DO_COUNT:.*]]
+}
+
+void do_disable()
+{
+// CHECK-LABEL: do_disable
+    int i = 1000;
+    __attribute__((opencl_unroll_hint(1)))
+    do {} while(i--> 0);
+// CHECK: br i1 %{{.*}}, label %{{.*}}, label %{{.*}}, !llvm.loop ![[DO_DISABLE:.*]]
+}
+
+void do_full()
+{
+// CHECK-LABEL: do_full
+    int i = 1000;
+    __attribute__((opencl_unroll_hint))
+    do {} while(i--> 0);
+// CHECK: br i1 %{{.*}}, label %{{.*}}, label %{{.*}}, !llvm.loop ![[DO_FULL:.*]]
+}
+
+
+// CHECK: ![[FOR_COUNT]]     =  distinct !{![[FOR_COUNT]],  ![[COUNT:.*]]}
+// CHECK: ![[COUNT]]         =  !{!"llvm.loop.unroll.count", i32 8}
+// CHECK: ![[FOR_DISABLE]]   =  distinct !{![[FOR_DISABLE]],  ![[DISABLE:.*]]}
+// CHECK: ![[DISABLE]]       =  !{!"llvm.loop.unroll.disable"}
+// CHECK: ![[FOR_FULL]]      =  distinct !{![[FOR_FULL]],  ![[FULL:.*]]}
+// CHECK: ![[FULL]]          =  !{!"llvm.loop.unroll.full"}
+// CHECK: ![[WHILE_COUNT]]   =  distinct !{![[WHILE_COUNT]],    ![[COUNT]]}
+// CHECK: ![[WHILE_DISABLE]] =  distinct !{![[WHILE_DISABLE]],  ![[DISABLE]]}
+// CHECK: ![[WHILE_FULL]]    =  distinct !{![[WHILE_FULL]],     ![[FULL]]}
+// CHECK: ![[DO_COUNT]]      =  distinct !{![[DO_COUNT]],       ![[COUNT]]}
+// CHECK: ![[DO_DISABLE]]    =  distinct !{![[DO_DISABLE]],     ![[DISABLE]]}
+// CHECK: ![[DO_FULL]]       =  distinct !{![[DO_FULL]],        ![[FULL]]}
diff --git a/test/CodeGenOpenCL/vla.cl b/test/CodeGenOpenCL/vla.cl
new file mode 100644
index 000000000000..cbf98445ce00
--- /dev/null
+++ b/test/CodeGenOpenCL/vla.cl
@@ -0,0 +1,18 @@
+// RUN: %clang_cc1 -emit-llvm -triple "spir-unknown-unknown" -O0 -cl-std=CL2.0 -o - %s | FileCheck %s
+
+constant int sz0 = 5;
+// CHECK: @sz0 = addrspace(2) constant i32 5
+const global int sz1 = 16;
+// CHECK: @sz1 = addrspace(1) constant i32 16
+const constant int sz2 = 8;
+// CHECK: @sz2 = addrspace(2) constant i32 8
+// CHECK: @testvla.vla2 = internal addrspace(3) global [8 x i16] undef
+
+kernel void testvla()
+{
+  int vla0[sz0];
+// CHECK: %vla0 = alloca [5 x i32]
+  char vla1[sz1];
+// CHECK: %vla1 = alloca [16 x i8]
+  local short vla2[sz2];
+}