summaryrefslogtreecommitdiff
path: root/clang/test/CodeGenOpenCL
diff options
context:
space:
mode:
Diffstat (limited to 'clang/test/CodeGenOpenCL')
-rw-r--r--clang/test/CodeGenOpenCL/2011-04-15-vec-init-from-vec.cl11
-rw-r--r--clang/test/CodeGenOpenCL/address-spaces.cl27
-rw-r--r--clang/test/CodeGenOpenCL/ext-vector-shuffle.cl17
-rw-r--r--clang/test/CodeGenOpenCL/fpmath.cl25
-rw-r--r--clang/test/CodeGenOpenCL/kernel-metadata.cl10
-rw-r--r--clang/test/CodeGenOpenCL/local.cl7
-rw-r--r--clang/test/CodeGenOpenCL/ptx-calls.cl12
-rw-r--r--clang/test/CodeGenOpenCL/ptx-kernels.cl10
-rw-r--r--clang/test/CodeGenOpenCL/single-precision-constant.cl7
-rw-r--r--clang/test/CodeGenOpenCL/vector_literals_nested.cl23
-rw-r--r--clang/test/CodeGenOpenCL/vector_literals_valid.cl22
-rw-r--r--clang/test/CodeGenOpenCL/vector_logops.cl19
12 files changed, 190 insertions, 0 deletions
diff --git a/clang/test/CodeGenOpenCL/2011-04-15-vec-init-from-vec.cl b/clang/test/CodeGenOpenCL/2011-04-15-vec-init-from-vec.cl
new file mode 100644
index 0000000..fbe3d89
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/2011-04-15-vec-init-from-vec.cl
@@ -0,0 +1,11 @@
+// RUN: %clang_cc1 %s -emit-llvm -o %t
+
+typedef __attribute__((ext_vector_type(4))) unsigned char uchar4;
+typedef __attribute__((ext_vector_type(8))) unsigned char uchar8;
+
+// OpenCL allows vectors to be initialized by vectors Handle bug in
+// VisitInitListExpr for this case below.
+void foo( uchar8 x )
+{
+ uchar4 val[4] = {{(uchar4){x.lo}}};
+}
diff --git a/clang/test/CodeGenOpenCL/address-spaces.cl b/clang/test/CodeGenOpenCL/address-spaces.cl
new file mode 100644
index 0000000..e030c77
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/address-spaces.cl
@@ -0,0 +1,27 @@
+// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s
+
+void f__p(__private int *arg) { }
+// CHECK: i32* nocapture %arg
+
+void f__g(__global int *arg) { }
+// CHECK: i32 addrspace(1)* nocapture %arg
+
+void f__l(__local int *arg) { }
+// CHECK: i32 addrspace(2)* nocapture %arg
+
+void f__c(__constant int *arg) { }
+// CHECK: i32 addrspace(3)* nocapture %arg
+
+
+void fp(private int *arg) { }
+// CHECK: i32* nocapture %arg
+
+void fg(global int *arg) { }
+// CHECK: i32 addrspace(1)* nocapture %arg
+
+void fl(local int *arg) { }
+// CHECK: i32 addrspace(2)* nocapture %arg
+
+void fc(constant int *arg) { }
+// CHECK: i32 addrspace(3)* nocapture %arg
+
diff --git a/clang/test/CodeGenOpenCL/ext-vector-shuffle.cl b/clang/test/CodeGenOpenCL/ext-vector-shuffle.cl
new file mode 100644
index 0000000..ee88ba3
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/ext-vector-shuffle.cl
@@ -0,0 +1,17 @@
+// RUN: %clang_cc1 %s -cl-opt-disable -emit-llvm -o - | not grep 'extractelement'
+// RUN: %clang_cc1 %s -cl-opt-disable -emit-llvm -o - | not grep 'insertelement'
+// RUN: %clang_cc1 %s -cl-opt-disable -emit-llvm -o - | grep 'shufflevector'
+
+typedef __attribute__(( ext_vector_type(2) )) float float2;
+typedef __attribute__(( ext_vector_type(4) )) float float4;
+
+float2 test1(float4 V) {
+ return V.xy + V.wz;
+}
+
+float4 test2(float4 V) {
+ float2 W = V.ww;
+ return W.xyxy + W.yxyx;
+}
+
+float4 test3(float4 V1, float4 V2) { return (float4)(V1.zw, V2.xy); }
diff --git a/clang/test/CodeGenOpenCL/fpmath.cl b/clang/test/CodeGenOpenCL/fpmath.cl
new file mode 100644
index 0000000..704fcd7
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/fpmath.cl
@@ -0,0 +1,25 @@
+// RUN: %clang_cc1 %s -emit-llvm -o - | FileCheck %s
+
+typedef __attribute__(( ext_vector_type(4) )) float float4;
+
+float spscalardiv(float a, float b) {
+ // CHECK: @spscalardiv
+ // CHECK: fdiv{{.*}}, !fpmath ![[MD:[0-9]+]]
+ return a / b;
+}
+
+float4 spvectordiv(float4 a, float4 b) {
+ // CHECK: @spvectordiv
+ // CHECK: fdiv{{.*}}, !fpmath ![[MD]]
+ return a / b;
+}
+
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+double dpscalardiv(double a, double b) {
+ // CHECK: @dpscalardiv
+ // CHECK-NOT: !fpmath
+ return a / b;
+}
+
+// CHECK: ![[MD]] = metadata !{float 2.500000e+00}
diff --git a/clang/test/CodeGenOpenCL/kernel-metadata.cl b/clang/test/CodeGenOpenCL/kernel-metadata.cl
new file mode 100644
index 0000000..3e10a11
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/kernel-metadata.cl
@@ -0,0 +1,10 @@
+// RUN: %clang_cc1 %s -emit-llvm -o - | FileCheck %s
+
+void normal_function() {
+}
+
+__kernel void kernel_function() {
+}
+
+// CHECK: !opencl.kernels = !{!0}
+// CHECK: !0 = metadata !{void ()* @kernel_function}
diff --git a/clang/test/CodeGenOpenCL/local.cl b/clang/test/CodeGenOpenCL/local.cl
new file mode 100644
index 0000000..32fa7be
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/local.cl
@@ -0,0 +1,7 @@
+// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s
+
+__kernel void foo(void) {
+ // CHECK: @foo.i = internal addrspace(2)
+ __local int i;
+ ++i;
+}
diff --git a/clang/test/CodeGenOpenCL/ptx-calls.cl b/clang/test/CodeGenOpenCL/ptx-calls.cl
new file mode 100644
index 0000000..6f33640
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/ptx-calls.cl
@@ -0,0 +1,12 @@
+// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
+
+void device_function() {
+}
+// CHECK: define ptx_device void @device_function()
+
+__kernel void kernel_function() {
+ device_function();
+}
+// CHECK: define ptx_kernel void @kernel_function()
+// CHECK: call ptx_device void @device_function()
+
diff --git a/clang/test/CodeGenOpenCL/ptx-kernels.cl b/clang/test/CodeGenOpenCL/ptx-kernels.cl
new file mode 100644
index 0000000..4d6fa10
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/ptx-kernels.cl
@@ -0,0 +1,10 @@
+// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -emit-llvm -o - | FileCheck %s
+
+void device_function() {
+}
+// CHECK: define ptx_device void @device_function()
+
+__kernel void kernel_function() {
+}
+// CHECK: define ptx_kernel void @kernel_function()
+
diff --git a/clang/test/CodeGenOpenCL/single-precision-constant.cl b/clang/test/CodeGenOpenCL/single-precision-constant.cl
new file mode 100644
index 0000000..62b37c1
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/single-precision-constant.cl
@@ -0,0 +1,7 @@
+// RUN: %clang_cc1 %s -cl-single-precision-constant -emit-llvm -o - | FileCheck %s
+
+float fn(float f) {
+ // CHECK: fmul float
+ // CHECK: fadd float
+ return f*2. + 1.;
+}
diff --git a/clang/test/CodeGenOpenCL/vector_literals_nested.cl b/clang/test/CodeGenOpenCL/vector_literals_nested.cl
new file mode 100644
index 0000000..b9013d0
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/vector_literals_nested.cl
@@ -0,0 +1,23 @@
+// RUN: %clang_cc1 %s -emit-llvm -O3 -o - | FileCheck %s
+
+typedef int int2 __attribute((ext_vector_type(2)));
+typedef int int4 __attribute((ext_vector_type(4)));
+
+__constant const int4 itest1 = (int4)(1, 2, ((int2)(3, 4)));
+// CHECK: constant <4 x i32> <i32 1, i32 2, i32 3, i32 4>
+__constant const int4 itest2 = (int4)(1, 2, ((int2)(3)));
+// CHECK: constant <4 x i32> <i32 1, i32 2, i32 3, i32 3>
+
+typedef float float2 __attribute((ext_vector_type(2)));
+typedef float float4 __attribute((ext_vector_type(4)));
+
+void ftest1(float4 *p) {
+ *p = (float4)(1.1f, 1.2f, ((float2)(1.3f, 1.4f)));
+// CHECK: store <4 x float> <float 0x3FF19999A0000000, float 0x3FF3333340000000, float 0x3FF4CCCCC0000000, float 0x3FF6666660000000>
+}
+
+float4 ftest2(float4 *p) {
+ *p = (float4)(1.1f, 1.2f, ((float2)(1.3f)));
+// CHECK: store <4 x float> <float 0x3FF19999A0000000, float 0x3FF3333340000000, float 0x3FF4CCCCC0000000, float 0x3FF4CCCCC0000000>
+}
+
diff --git a/clang/test/CodeGenOpenCL/vector_literals_valid.cl b/clang/test/CodeGenOpenCL/vector_literals_valid.cl
new file mode 100644
index 0000000..bba5b23
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/vector_literals_valid.cl
@@ -0,0 +1,22 @@
+// RUN: %clang_cc1 -emit-llvm %s -o %t
+
+typedef __attribute__(( ext_vector_type(2) )) int int2;
+typedef __attribute__(( ext_vector_type(3) )) int int3;
+typedef __attribute__(( ext_vector_type(4) )) int int4;
+typedef __attribute__(( ext_vector_type(8) )) int int8;
+typedef __attribute__(( ext_vector_type(4) )) float float4;
+
+void vector_literals_valid() {
+ int4 a_1_1_1_1 = (int4)(1,2,3,4);
+ int4 a_2_1_1 = (int4)((int2)(1,2),3,4);
+ int4 a_1_2_1 = (int4)(1,(int2)(2,3),4);
+ int4 a_1_1_2 = (int4)(1,2,(int2)(3,4));
+ int4 a_2_2 = (int4)((int2)(1,2),(int2)(3,4));
+ int4 a_3_1 = (int4)((int3)(1,2,3),4);
+ int4 a_1_3 = (int4)(1,(int3)(2,3,4));
+ int4 a = (int4)(1);
+ int8 b = (int8)(1,2,a.xy,a);
+ float4 V2 = (float4) (1);
+}
+
+
diff --git a/clang/test/CodeGenOpenCL/vector_logops.cl b/clang/test/CodeGenOpenCL/vector_logops.cl
new file mode 100644
index 0000000..388f1d7
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/vector_logops.cl
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 -O3 %s -emit-llvm -o - | FileCheck %s
+
+typedef int int2 __attribute((ext_vector_type(2)));
+
+int test1()
+{
+ int2 a = (int2)(1,0);
+ int2 b = (int2)(1,1);
+ return (a&&b).x + (a||b).y;
+ // CHECK: ret i32 -2
+}
+
+int test2()
+{
+ int2 a = (int2)(1,0);
+ return (!a).y;
+ // CHECK: ret i32 -1
+}
+
generated by cgit v1.2.3 (git 2.46.2) at 2024-11-10 07:06:39 +0000