7 files changed, 73 insertions, 8 deletions
diff --git a/test/CodeGenOpenCL/bool_cast.cl b/test/CodeGenOpenCL/bool_cast.cl
index ab40eccf571f7..72926eb141ab0 100644
--- a/test/CodeGenOpenCL/bool_cast.cl
+++ b/test/CodeGenOpenCL/bool_cast.cl
@@ -4,7 +4,7 @@ typedef unsigned char uchar4 __attribute((ext_vector_type(4)));
 typedef unsigned int int4 __attribute((ext_vector_type(4)));
 typedef float float4 __attribute((ext_vector_type(4)));
 
-// CHECK-LABEL: define void @ker()
+// CHECK-LABEL: define spir_kernel void @ker()
 void kernel ker() {
   bool t = true;
   int4 vec4 = (int4)t;
diff --git a/test/CodeGenOpenCL/kernel-attributes.cl b/test/CodeGenOpenCL/kernel-attributes.cl
index e61a75f7b5370..eff2e57d8c9b6 100644
--- a/test/CodeGenOpenCL/kernel-attributes.cl
+++ b/test/CodeGenOpenCL/kernel-attributes.cl
@@ -3,13 +3,13 @@
 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]+]]
+// CHECK: define spir_kernel 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: define spir_kernel void @kernel2(i32 {{[^%]*}}%a) {{[^{]+}} !vec_type_hint ![[MD3:[0-9]+]] !work_group_size_hint ![[MD4:[0-9]+]]
 
 kernel __attribute__((intel_reqd_sub_group_size(8))) void kernel3(int a) {}
-// CHECK: define void @kernel3(i32 {{[^%]*}}%a) {{[^{]+}} !intel_reqd_sub_group_size ![[MD5:[0-9]+]]
+// CHECK: define spir_kernel void @kernel3(i32 {{[^%]*}}%a) {{[^{]+}} !intel_reqd_sub_group_size ![[MD5:[0-9]+]]
 
 // CHECK: [[MD1]] = !{i32 undef, i32 1}
 // CHECK: [[MD2]] = !{i32 1, i32 2, i32 4}
diff --git a/test/CodeGenOpenCL/kernel-metadata.cl b/test/CodeGenOpenCL/kernel-metadata.cl
index 4165f1fa0ce58..95be43015aa96 100644
--- a/test/CodeGenOpenCL/kernel-metadata.cl
+++ b/test/CodeGenOpenCL/kernel-metadata.cl
@@ -6,5 +6,5 @@ void normal_function() {
 __kernel void kernel_function() {
 }
 
-// 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: define spir_kernel 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/kernels-have-spir-cc-by-default.cl b/test/CodeGenOpenCL/kernels-have-spir-cc-by-default.cl
new file mode 100644
index 0000000000000..5bb52e9beb514
--- /dev/null
+++ b/test/CodeGenOpenCL/kernels-have-spir-cc-by-default.cl
@@ -0,0 +1,65 @@
+// RUN: %clang_cc1 %s -cl-std=CL1.2 -emit-llvm -triple x86_64-unknown-unknown -o - | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=CL1.2 -emit-llvm -triple amdgcn-unknown-unknown -o - | FileCheck -check-prefixes=AMDGCN %s
+// Test that the kernels always use the SPIR calling convention
+// to have unambiguous mapping of arguments to feasibly implement
+// clSetKernelArg().
+
+typedef struct int_single {
+    int a;
+} int_single;
+
+typedef struct int_pair {
+    long a;
+    long b;
+} int_pair;
+
+typedef struct test_struct {
+    int elementA;
+    int elementB;
+    long elementC;
+    char elementD;
+    long elementE;
+    float elementF;
+    short elementG;
+    double elementH;
+} test_struct;
+
+kernel void test_single(int_single input, global int* output) {
+// CHECK: spir_kernel
+// AMDGCN: define amdgpu_kernel void @test_single
+// CHECK: struct.int_single* byval nocapture
+// CHECK: i32* nocapture %output
+ output[0] = input.a;
+}
+
+kernel void test_pair(int_pair input, global int* output) {
+// CHECK: spir_kernel
+// AMDGCN: define amdgpu_kernel void @test_pair
+// CHECK: struct.int_pair* byval nocapture
+// CHECK: i32* nocapture %output
+ output[0] = (int)input.a;
+ output[1] = (int)input.b;
+}
+
+kernel void test_kernel(test_struct input, global int* output) {
+// CHECK: spir_kernel
+// AMDGCN: define amdgpu_kernel void @test_kernel
+// CHECK: struct.test_struct* byval nocapture
+// CHECK: i32* nocapture %output
+ output[0] = input.elementA;
+ output[1] = input.elementB;
+ output[2] = (int)input.elementC;
+ output[3] = (int)input.elementD;
+ output[4] = (int)input.elementE;
+ output[5] = (int)input.elementF;
+ output[6] = (int)input.elementG;
+ output[7] = (int)input.elementH;
+};
+
+void test_function(int_pair input, global int* output) {
+// CHECK-NOT: spir_kernel
+// AMDGCN-NOT: define amdgpu_kernel void @test_function
+// CHECK: i64 %input.coerce0, i64 %input.coerce1, i32* nocapture %output
+ output[0] = (int)input.a;
+ output[1] = (int)input.b;
+}
diff --git a/test/CodeGenOpenCL/pipe_types.cl b/test/CodeGenOpenCL/pipe_types.cl
index f5b42e2a490fc..7c11f74ad7b96 100644
--- a/test/CodeGenOpenCL/pipe_types.cl
+++ b/test/CodeGenOpenCL/pipe_types.cl
@@ -28,7 +28,7 @@ void test5(read_only pipe int4 p) {
 
 typedef read_only pipe int MyPipe;
 kernel void test6(MyPipe p) {
-// CHECK: define void @test6(%opencl.pipe_t* %p)
+// CHECK: define spir_kernel void @test6(%opencl.pipe_t* %p)
 }
 
 struct Person {
diff --git a/test/CodeGenOpenCL/ptx-calls.cl b/test/CodeGenOpenCL/ptx-calls.cl
index bde00bc3d73a7..2a3400371edf3 100644
--- a/test/CodeGenOpenCL/ptx-calls.cl
+++ b/test/CodeGenOpenCL/ptx-calls.cl
@@ -7,7 +7,7 @@ void device_function() {
 __kernel void kernel_function() {
   device_function();
 }
-// CHECK-LABEL: define void @kernel_function()
+// CHECK-LABEL: define spir_kernel void @kernel_function()
 // CHECK: call void @device_function()
 // CHECK: !{{[0-9]+}} = !{void ()* @kernel_function, !"kernel", i32 1}
 
diff --git a/test/CodeGenOpenCL/ptx-kernels.cl b/test/CodeGenOpenCL/ptx-kernels.cl
index fc6de4f3d517a..b9e1c224c7ce2 100644
--- a/test/CodeGenOpenCL/ptx-kernels.cl
+++ b/test/CodeGenOpenCL/ptx-kernels.cl
@@ -6,6 +6,6 @@ void device_function() {
 
 __kernel void kernel_function() {
 }
-// CHECK-LABEL: define void @kernel_function()
+// CHECK-LABEL: define spir_kernel void @kernel_function()
 
 // CHECK: !{{[0-9]+}} = !{void ()* @kernel_function, !"kernel", i32 1}