summaryrefslogtreecommitdiffstats
path: root/test/CodeGenOpenCL
diff options
context:
space:
mode:
Diffstat (limited to 'test/CodeGenOpenCL')
-rw-r--r--test/CodeGenOpenCL/addr-space-struct-arg.cl23
-rw-r--r--test/CodeGenOpenCL/event_t.cl12
-rw-r--r--test/CodeGenOpenCL/half.cl15
-rw-r--r--test/CodeGenOpenCL/kernel-arg-info.cl19
-rw-r--r--test/CodeGenOpenCL/kernel-attributes.cl12
-rw-r--r--test/CodeGenOpenCL/local.cl5
-rw-r--r--test/CodeGenOpenCL/logical-ops.cl56
-rw-r--r--test/CodeGenOpenCL/opencl_types.cl37
-rw-r--r--test/CodeGenOpenCL/ptx-calls.cl7
-rw-r--r--test/CodeGenOpenCL/ptx-kernels.cl5
-rw-r--r--test/CodeGenOpenCL/shifts.cl57
-rw-r--r--test/CodeGenOpenCL/spir32_target.cl22
-rw-r--r--test/CodeGenOpenCL/spir64_target.cl21
13 files changed, 279 insertions, 12 deletions
diff --git a/test/CodeGenOpenCL/addr-space-struct-arg.cl b/test/CodeGenOpenCL/addr-space-struct-arg.cl
new file mode 100644
index 0000000..f04923d
--- /dev/null
+++ b/test/CodeGenOpenCL/addr-space-struct-arg.cl
@@ -0,0 +1,23 @@
+// RUN: %clang_cc1 %s -emit-llvm -o - -O0 -ffake-address-space-map -triple i686-pc-darwin | FileCheck %s
+
+typedef struct {
+ int cells[9];
+} Mat3X3;
+
+typedef struct {
+ int cells[16];
+} Mat4X4;
+
+Mat4X4 __attribute__((noinline)) foo(Mat3X3 in) {
+ Mat4X4 out;
+ return out;
+}
+
+kernel void ker(global Mat3X3 *in, global Mat4X4 *out) {
+ out[0] = foo(in[1]);
+}
+
+// Expect two mem copies: one for the argument "in", and one for
+// the return value.
+// CHECK: call void @llvm.memcpy.p0i8.p1i8.i32(i8*
+// CHECK: call void @llvm.memcpy.p1i8.p0i8.i32(i8 addrspace(1)*
diff --git a/test/CodeGenOpenCL/event_t.cl b/test/CodeGenOpenCL/event_t.cl
new file mode 100644
index 0000000..ddf12a9
--- /dev/null
+++ b/test/CodeGenOpenCL/event_t.cl
@@ -0,0 +1,12 @@
+// RUN: %clang_cc1 %s -emit-llvm -o - -O0 | FileCheck %s
+
+void foo(event_t evt);
+
+void kernel ker() {
+ event_t e;
+// CHECK: alloca %opencl.event_t*,
+ foo(e);
+// CHECK: call void @foo(%opencl.event_t* %
+ foo(0);
+// CHECK: call void @foo(%opencl.event_t* null)
+}
diff --git a/test/CodeGenOpenCL/half.cl b/test/CodeGenOpenCL/half.cl
new file mode 100644
index 0000000..7ecae89
--- /dev/null
+++ b/test/CodeGenOpenCL/half.cl
@@ -0,0 +1,15 @@
+// RUN: %clang_cc1 %s -emit-llvm -o - | FileCheck %s
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+
+half test()
+{
+ half x = 0.1f;
+ x+=2.0f;
+ x-=2.0f;
+ half y = x + x;
+ half z = y * 1.0f;
+ return z;
+// CHECK: half 0xH3260
+}
diff --git a/test/CodeGenOpenCL/kernel-arg-info.cl b/test/CodeGenOpenCL/kernel-arg-info.cl
index 9d52736..c7e2049 100644
--- a/test/CodeGenOpenCL/kernel-arg-info.cl
+++ b/test/CodeGenOpenCL/kernel-arg-info.cl
@@ -1,7 +1,20 @@
-// RUN: %clang_cc1 %s -cl-kernel-arg-info -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -cl-kernel-arg-info -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
-kernel void foo(int *X, int Y, int anotherArg) {
+kernel void foo(__global int * restrict X, const int Y,
+ volatile int anotherArg, __constant float * restrict Z) {
*X = Y + anotherArg;
}
-// CHECK: metadata !{metadata !"kernel_arg_name", metadata !"X", metadata !"Y", metadata !"anotherArg"}
+// CHECK: metadata !{metadata !"kernel_arg_addr_space", i32 1, i32 0, i32 0, i32 2}
+// CHECK: metadata !{metadata !"kernel_arg_access_qual", metadata !"none", metadata !"none", metadata !"none", metadata !"none"}
+// CHECK: metadata !{metadata !"kernel_arg_type", metadata !"int*", metadata !"int", metadata !"int", metadata !"float*"}
+// CHECK: metadata !{metadata !"kernel_arg_type_qual", metadata !"restrict", metadata !"const", metadata !"volatile", metadata !"restrict const"}
+// CHECK: metadata !{metadata !"kernel_arg_name", metadata !"X", metadata !"Y", metadata !"anotherArg", metadata !"Z"}
+
+kernel void foo2(read_only image1d_t img1, image2d_t img2, write_only image2d_array_t img3) {
+}
+// CHECK: metadata !{metadata !"kernel_arg_addr_space", i32 0, i32 0, i32 0}
+// CHECK: metadata !{metadata !"kernel_arg_access_qual", metadata !"read_only", metadata !"read_only", metadata !"write_only"}
+// CHECK: metadata !{metadata !"kernel_arg_type", metadata !"image1d_t", metadata !"image2d_t", metadata !"image2d_array_t"}
+// CHECK: metadata !{metadata !"kernel_arg_type_qual", metadata !"", metadata !"", metadata !""}
+// CHECK: metadata !{metadata !"kernel_arg_name", metadata !"img1", metadata !"img2", metadata !"img3"}
diff --git a/test/CodeGenOpenCL/kernel-attributes.cl b/test/CodeGenOpenCL/kernel-attributes.cl
index de16a41..1166f93 100644
--- a/test/CodeGenOpenCL/kernel-attributes.cl
+++ b/test/CodeGenOpenCL/kernel-attributes.cl
@@ -1,12 +1,16 @@
// RUN: %clang_cc1 -emit-llvm -O0 -o - %s | FileCheck %s
-kernel __attribute__((reqd_work_group_size(1,2,4))) void kernel1(int a) {}
+typedef unsigned int uint4 __attribute__((ext_vector_type(4)));
-kernel __attribute__((work_group_size_hint(8,16,32))) void kernel2(int a) {}
+kernel __attribute__((vec_type_hint(int))) __attribute__((reqd_work_group_size(1,2,4))) void kernel1(int a) {}
+
+kernel __attribute__((vec_type_hint(uint4))) __attribute__((work_group_size_hint(8,16,32))) void kernel2(int a) {}
// CHECK: opencl.kernels = !{[[MDNODE0:![0-9]+]], [[MDNODE3:![0-9]+]]}
-// CHECK: [[MDNODE0]] = metadata !{void (i32)* @kernel1, metadata [[MDNODE2:![0-9]+]]}
+// CHECK: [[MDNODE0]] = metadata !{void (i32)* @kernel1, metadata [[MDNODE1:![0-9]+]], metadata [[MDNODE2:![0-9]+]]}
+// CHECK: [[MDNODE1]] = metadata !{metadata !"vec_type_hint", i32 undef, i32 1}
// CHECK: [[MDNODE2]] = metadata !{metadata !"reqd_work_group_size", i32 1, i32 2, i32 4}
-// CHECK: [[MDNODE3]] = metadata !{void (i32)* @kernel2, metadata [[MDNODE5:![0-9]+]]}
+// CHECK: [[MDNODE3]] = metadata !{void (i32)* @kernel2, metadata [[MDNODE4:![0-9]+]], metadata [[MDNODE5:![0-9]+]]}
+// CHECK: [[MDNODE4]] = metadata !{metadata !"vec_type_hint", <4 x i32> undef, i32 0}
// CHECK: [[MDNODE5]] = metadata !{metadata !"work_group_size_hint", i32 8, i32 16, i32 32}
diff --git a/test/CodeGenOpenCL/local.cl b/test/CodeGenOpenCL/local.cl
index 32fa7be..b4bd008 100644
--- a/test/CodeGenOpenCL/local.cl
+++ b/test/CodeGenOpenCL/local.cl
@@ -5,3 +5,8 @@ __kernel void foo(void) {
__local int i;
++i;
}
+
+// CHECK: define void @_Z3barPU3AS2i
+__kernel void __attribute__((__overloadable__)) bar(local int *x) {
+ *x = 5;
+}
diff --git a/test/CodeGenOpenCL/logical-ops.cl b/test/CodeGenOpenCL/logical-ops.cl
new file mode 100644
index 0000000..ac1c1b5
--- /dev/null
+++ b/test/CodeGenOpenCL/logical-ops.cl
@@ -0,0 +1,56 @@
+// RUN: %clang_cc1 %s -emit-llvm -o - -cl-std=CL1.2 -O1 -triple x86_64-unknown-linux-gnu | FileCheck %s
+
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+typedef int int4 __attribute((ext_vector_type(4)));
+typedef long long4 __attribute((ext_vector_type(4)));
+typedef float float4 __attribute((ext_vector_type(4)));
+typedef double double4 __attribute((ext_vector_type(4)));
+
+// CHECK: floatops
+kernel void floatops(global int4 *out, global float4 *fout) {
+ // CHECK: store <4 x i32> <i32 -1, i32 -1, i32 -1, i32 -1>
+ out[0] = (float4)(1, 1, 1, 1) && 1.0f;
+ // CHECK: store <4 x i32> zeroinitializer
+ out[1] = (float4)(0, 0, 0, 0) && (float4)(0, 0, 0, 0);
+
+ // CHECK: store <4 x i32> <i32 -1, i32 -1, i32 -1, i32 -1>
+ out[2] = (float4)(0, 0, 0, 0) || (float4)(1, 1, 1, 1);
+ // CHECK: store <4 x i32> zeroinitializer
+ out[3] = (float4)(0, 0, 0, 0) || 0.0f;
+
+ // CHECK: store <4 x i32> <i32 -1, i32 -1, i32 -1, i32 -1>
+ out[4] = !(float4)(0, 0, 0, 0);
+ // CHECK: store <4 x i32> zeroinitializer
+ out[5] = !(float4)(1, 2, 3, 4);
+ // CHECK: store <4 x i32> <i32 -1, i32 0, i32 -1, i32 0>
+ out[6] = !(float4)(0, 1, 0, 1);
+ // CHECK: store <4 x float> <float 1.000000e+00, float 1.000000e+00, float 1.000000e+00, float 1.000000e+00>
+ fout[0] = (float4)(!0.0f);
+ // CHECK: store <4 x float> zeroinitializer
+ fout[1] = (float4)(!1.0f);
+}
+
+// CHECK: doubleops
+kernel void doubleops(global long4 *out, global double4 *dout) {
+ // CHECK: store <4 x i64> <i64 -1, i64 -1, i64 -1, i64 -1>
+ out[0] = (double4)(1, 1, 1, 1) && 1.0;
+ // CHECK: store <4 x i64> zeroinitializer
+ out[1] = (double4)(0, 0, 0, 0) && (double4)(0, 0, 0, 0);
+
+ // CHECK: store <4 x i64> <i64 -1, i64 -1, i64 -1, i64 -1>
+ out[2] = (double4)(0, 0, 0, 0) || (double4)(1, 1, 1, 1);
+ // CHECK: store <4 x i64> zeroinitializer
+ out[3] = (double4)(0, 0, 0, 0) || 0.0f;
+
+ // CHECK: store <4 x i64> <i64 -1, i64 -1, i64 -1, i64 -1>
+ out[4] = !(double4)(0, 0, 0, 0);
+ // CHECK: store <4 x i64> zeroinitializer
+ out[5] = !(double4)(1, 2, 3, 4);
+ // CHECK: store <4 x i64> <i64 -1, i64 0, i64 -1, i64 0>
+ out[6] = !(double4)(0, 1, 0, 1);
+ // CHECK: store <4 x double> <double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00>
+ dout[0] = (double4)(!0.0f);
+ // CHECK: store <4 x double> zeroinitializer
+ dout[1] = (double4)(!1.0f);
+}
diff --git a/test/CodeGenOpenCL/opencl_types.cl b/test/CodeGenOpenCL/opencl_types.cl
new file mode 100644
index 0000000..b1e558d
--- /dev/null
+++ b/test/CodeGenOpenCL/opencl_types.cl
@@ -0,0 +1,37 @@
+// RUN: %clang_cc1 %s -emit-llvm -o - -O0 | FileCheck %s
+
+constant sampler_t glb_smp = 7;
+// CHECK: global i32 7
+
+void fnc1(image1d_t img) {}
+// CHECK: @fnc1(%opencl.image1d_t*
+
+void fnc1arr(image1d_array_t img) {}
+// CHECK: @fnc1arr(%opencl.image1d_array_t*
+
+void fnc1buff(image1d_buffer_t img) {}
+// CHECK: @fnc1buff(%opencl.image1d_buffer_t*
+
+void fnc2(image2d_t img) {}
+// CHECK: @fnc2(%opencl.image2d_t*
+
+void fnc2arr(image2d_array_t img) {}
+// CHECK: @fnc2arr(%opencl.image2d_array_t*
+
+void fnc3(image3d_t img) {}
+// CHECK: @fnc3(%opencl.image3d_t*
+
+void fnc4smp(sampler_t s) {}
+// CHECK: 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,
+ fnc4smp(smp);
+// CHECK: call void @fnc4smp(i32
+ fnc4smp(glb_smp);
+// CHECK: call void @fnc4smp(i32
+}
diff --git a/test/CodeGenOpenCL/ptx-calls.cl b/test/CodeGenOpenCL/ptx-calls.cl
index 34a21c6..d990451 100644
--- a/test/CodeGenOpenCL/ptx-calls.cl
+++ b/test/CodeGenOpenCL/ptx-calls.cl
@@ -2,11 +2,12 @@
void device_function() {
}
-// CHECK: define ptx_device void @device_function()
+// CHECK: define void @device_function()
__kernel void kernel_function() {
device_function();
}
-// CHECK: define ptx_kernel void @kernel_function()
-// CHECK: call ptx_device void @device_function()
+// CHECK: define void @kernel_function()
+// CHECK: call void @device_function()
+// CHECK: !{{[0-9]+}} = metadata !{void ()* @kernel_function, metadata !"kernel", i32 1}
diff --git a/test/CodeGenOpenCL/ptx-kernels.cl b/test/CodeGenOpenCL/ptx-kernels.cl
index 1d7e497..07648e4 100644
--- a/test/CodeGenOpenCL/ptx-kernels.cl
+++ b/test/CodeGenOpenCL/ptx-kernels.cl
@@ -2,9 +2,10 @@
void device_function() {
}
-// CHECK: define ptx_device void @device_function()
+// CHECK: define void @device_function()
__kernel void kernel_function() {
}
-// CHECK: define ptx_kernel void @kernel_function()
+// CHECK: define void @kernel_function()
+// CHECK: !{{[0-9]+}} = metadata !{void ()* @kernel_function, metadata !"kernel", i32 1}
diff --git a/test/CodeGenOpenCL/shifts.cl b/test/CodeGenOpenCL/shifts.cl
new file mode 100644
index 0000000..015a777
--- /dev/null
+++ b/test/CodeGenOpenCL/shifts.cl
@@ -0,0 +1,57 @@
+// RUN: %clang_cc1 -x cl -O1 -emit-llvm %s -o - -triple x86_64-linux-gnu | FileCheck %s
+// OpenCL essentially reduces all shift amounts to the last word-size bits before evaluating.
+// Test this both for variables and constants evaluated in the front-end.
+
+
+//CHECK: @positiveShift32
+int positiveShift32(int a,int b) {
+ //CHECK: [[M32:%.+]] = and i32 %b, 31
+ //CHECK-NEXT: [[C32:%.+]] = shl i32 %a, [[M32]]
+ int c = a<<b;
+ int d = ((int)1)<<33;
+ //CHECK-NEXT: [[E32:%.+]] = add nsw i32 [[C32]], 2
+ int e = c + d;
+ //CHECK-NEXT: ret i32 [[E32]]
+ return e;
+}
+
+//CHECK: @positiveShift64
+long positiveShift64(long a,long b) {
+ //CHECK: [[M64:%.+]] = and i64 %b, 63
+ //CHECK-NEXT: [[C64:%.+]] = ashr i64 %a, [[M64]]
+ long c = a>>b;
+ long d = ((long)8)>>65;
+ //CHECK-NEXT: [[E64:%.+]] = add nsw i64 [[C64]], 4
+ long e = c + d;
+ //CHECK-NEXT: ret i64 [[E64]]
+ return e;
+}
+
+typedef __attribute__((ext_vector_type(4))) int int4;
+
+//CHECK: @vectorVectorTest
+int4 vectorVectorTest(int4 a,int4 b) {
+ //CHECK: [[VM:%.+]] = and <4 x i32> %b, <i32 31, i32 31, i32 31, i32 31>
+ //CHECK-NEXT: [[VC:%.+]] = shl <4 x i32> %a, [[VM]]
+ int4 c = a << b;
+ //CHECK-NEXT: [[VF:%.+]] = add <4 x i32> [[VC]], <i32 2, i32 4, i32 16, i32 8>
+ int4 d = {1, 1, 1, 1};
+ int4 e = {33, 34, -28, -29};
+ int4 f = c + (d << e);
+ //CHECK-NEXT: ret <4 x i32> [[VF]]
+ return f;
+}
+
+//CHECK: @vectorScalarTest
+int4 vectorScalarTest(int4 a,int b) {
+ //CHECK: [[SP0:%.+]] = insertelement <4 x i32> undef, i32 %b, i32 0
+ //CHECK: [[SP1:%.+]] = shufflevector <4 x i32> [[SP0]], <4 x i32> undef, <4 x i32> zeroinitializer
+ //CHECK: [[VSM:%.+]] = and <4 x i32> [[SP1]], <i32 31, i32 31, i32 31, i32 31>
+ //CHECK-NEXT: [[VSC:%.+]] = shl <4 x i32> %a, [[VSM]]
+ int4 c = a << b;
+ //CHECK-NEXT: [[VSF:%.+]] = add <4 x i32> [[VSC]], <i32 4, i32 4, i32 4, i32 4>
+ int4 d = {1, 1, 1, 1};
+ int4 f = c + (d << 34);
+ //CHECK-NEXT: ret <4 x i32> [[VSF]]
+ return f;
+}
diff --git a/test/CodeGenOpenCL/spir32_target.cl b/test/CodeGenOpenCL/spir32_target.cl
new file mode 100644
index 0000000..8f395b3
--- /dev/null
+++ b/test/CodeGenOpenCL/spir32_target.cl
@@ -0,0 +1,22 @@
+// RUN: %clang_cc1 %s -triple "spir-unknown-unknown" -emit-llvm -o - | FileCheck %s
+
+// CHECK: target triple = "spir-unknown-unknown"
+
+typedef struct {
+ char c;
+ void *v;
+ void *v2;
+} my_st;
+
+kernel void foo(global long *arg) {
+ int res1[sizeof(my_st) == 12 ? 1 : -1];
+ int res2[sizeof(void *) == 4 ? 1 : -1];
+ int res3[sizeof(arg) == 4 ? 1 : -1];
+
+ my_st *tmp = 0;
+
+ arg[0] = (long)(&tmp->v);
+//CHECK: store i64 4, i64 addrspace(1)*
+ arg[1] = (long)(&tmp->v2);
+//CHECK: store i64 8, i64 addrspace(1)*
+}
diff --git a/test/CodeGenOpenCL/spir64_target.cl b/test/CodeGenOpenCL/spir64_target.cl
new file mode 100644
index 0000000..245cd80
--- /dev/null
+++ b/test/CodeGenOpenCL/spir64_target.cl
@@ -0,0 +1,21 @@
+// RUN: %clang_cc1 %s -triple "spir64-unknown-unknown" -emit-llvm -o - | FileCheck %s
+
+// CHECK: target triple = "spir64-unknown-unknown"
+
+typedef struct {
+ char c;
+ void *v;
+ void *v2;
+} my_st;
+
+kernel void foo(global long *arg) {
+ int res1[sizeof(my_st) == 24 ? 1 : -1];
+ int res2[sizeof(void *) == 8 ? 1 : -1];
+ int res3[sizeof(arg) == 8 ? 1 : -1];
+
+ my_st *tmp = 0;
+ arg[3] = (long)(&tmp->v);
+//CHECK: store i64 8, i64 addrspace(1)*
+ arg[4] = (long)(&tmp->v2);
+//CHECK: store i64 16, i64 addrspace(1)*
+}
OpenPOWER on IntegriCloud