aboutsummaryrefslogtreecommitdiff
path: root/test/CodeGenOpenCL
diff options
context:
space:
mode:
Diffstat (limited to 'test/CodeGenOpenCL')
-rw-r--r--test/CodeGenOpenCL/addr-space-struct-arg.cl46
-rw-r--r--test/CodeGenOpenCL/address-space-constant-initializers.cl2
-rw-r--r--test/CodeGenOpenCL/address-spaces-conversions.cl22
-rw-r--r--test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl48
-rw-r--r--test/CodeGenOpenCL/builtins-r600.cl37
-rw-r--r--test/CodeGenOpenCL/const-str-array-decay.cl11
-rw-r--r--test/CodeGenOpenCL/constant-addr-space-globals.cl8
-rw-r--r--test/CodeGenOpenCL/denorms-are-zero.cl5
-rw-r--r--test/CodeGenOpenCL/fpmath.cl2
-rw-r--r--test/CodeGenOpenCL/kernel-arg-info.cl59
-rw-r--r--test/CodeGenOpenCL/kernel-attributes.cl12
-rw-r--r--test/CodeGenOpenCL/kernel-metadata.cl7
-rw-r--r--test/CodeGenOpenCL/local-initializer-undef.cl24
-rw-r--r--test/CodeGenOpenCL/local.cl6
-rw-r--r--test/CodeGenOpenCL/opencl_types.cl5
-rw-r--r--test/CodeGenOpenCL/ptx-calls.cl2
-rw-r--r--test/CodeGenOpenCL/ptx-kernels.cl2
-rw-r--r--test/CodeGenOpenCL/relaxed-fpmath.cl36
-rw-r--r--test/CodeGenOpenCL/str_literals.cl4
19 files changed, 279 insertions, 59 deletions
diff --git a/test/CodeGenOpenCL/addr-space-struct-arg.cl b/test/CodeGenOpenCL/addr-space-struct-arg.cl
index f04923d39c58..d711f78d4ef0 100644
--- a/test/CodeGenOpenCL/addr-space-struct-arg.cl
+++ b/test/CodeGenOpenCL/addr-space-struct-arg.cl
@@ -1,23 +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)*
+// 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/address-space-constant-initializers.cl b/test/CodeGenOpenCL/address-space-constant-initializers.cl
index ae8cedc1ca5b..079b0706f498 100644
--- a/test/CodeGenOpenCL/address-space-constant-initializers.cl
+++ b/test/CodeGenOpenCL/address-space-constant-initializers.cl
@@ -12,7 +12,7 @@ typedef struct {
} ConstantArrayPointerStruct;
// CHECK: %struct.ConstantArrayPointerStruct = type { float addrspace(3)* }
-// CHECK: addrspace(3) global %struct.ConstantArrayPointerStruct { float addrspace(3)* bitcast (i8 addrspace(3)* getelementptr (i8 addrspace(3)* bitcast (%struct.ArrayStruct addrspace(3)* @constant_array_struct to i8 addrspace(3)*), i64 4) to float addrspace(3)*) }
+// CHECK: addrspace(3) constant %struct.ConstantArrayPointerStruct { float addrspace(3)* bitcast (i8 addrspace(3)* getelementptr (i8 addrspace(3)* bitcast (%struct.ArrayStruct addrspace(3)* @constant_array_struct to i8 addrspace(3)*), i64 4) to float addrspace(3)*) }
// Bug 18567
__constant ConstantArrayPointerStruct constant_array_pointer_struct = {
&constant_array_struct.f
diff --git a/test/CodeGenOpenCL/address-spaces-conversions.cl b/test/CodeGenOpenCL/address-spaces-conversions.cl
new file mode 100644
index 000000000000..bc80f47f5f42
--- /dev/null
+++ b/test/CodeGenOpenCL/address-spaces-conversions.cl
@@ -0,0 +1,22 @@
+// RUN: %clang_cc1 %s -triple x86_64-unknown-linux-gnu -O0 -ffake-address-space-map -cl-std=CL2.0 -emit-llvm -o - | FileCheck %s
+
+// test that we generate address space casts everywhere we need conversions of
+// pointers to different address spaces
+
+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)*
+ arg_gen = &var_priv; // implicit cast with obtaining adr, private -> generic
+ // CHECK: %{{[0-9]+}} = addrspacecast i32* %var_priv to i32 addrspace(4)*
+ arg_glob = (global int *)arg_gen; // explicit cast
+ // CHECK: %{{[0-9]+}} = addrspacecast i32 addrspace(4)* %{{[0-9]+}} to i32 addrspace(1)*
+ 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)*
+ var_priv = arg_gen - arg_glob; // arithmetic operation
+ // CHECK: %{{.*}} = ptrtoint i32 addrspace(4)* %{{.*}} to i64
+ // CHECK: %{{.*}} = ptrtoint i32 addrspace(1)* %{{.*}} to i64
+ var_priv = arg_gen > arg_glob; // comparison
+ // CHECK: %{{[0-9]+}} = addrspacecast i32 addrspace(1)* %{{[0-9]+}} to i32 addrspace(4)*
+}
diff --git a/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl b/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl
new file mode 100644
index 000000000000..35bdcead3128
--- /dev/null
+++ b/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl
@@ -0,0 +1,48 @@
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -O0 -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86 %s
+
+// Make sure this is silently accepted on other targets.
+
+__attribute__((amdgpu_num_vgpr(64))) // expected-no-diagnostics
+kernel void test_num_vgpr64() {
+// CHECK: define 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]+]]
+}
+
+__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]+]]
+
+}
+
+__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]+]]
+}
+
+__attribute__((amdgpu_num_vgpr(0))) // expected-no-diagnostics
+kernel void test_num_vgpr0() {
+}
+
+__attribute__((amdgpu_num_sgpr(0))) // expected-no-diagnostics
+kernel void test_num_sgpr0() {
+}
+
+__attribute__((amdgpu_num_vgpr(0), amdgpu_num_sgpr(0))) // expected-no-diagnostics
+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-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"
+// CHECK-DAG: attributes [[ATTR_SGPR20_VGPR40]] = { nounwind "amdgpu_num_sgpr"="20" "amdgpu_num_vgpr"="40"
diff --git a/test/CodeGenOpenCL/builtins-r600.cl b/test/CodeGenOpenCL/builtins-r600.cl
index cebeba185ea3..3e416b0323c2 100644
--- a/test/CodeGenOpenCL/builtins-r600.cl
+++ b/test/CodeGenOpenCL/builtins-r600.cl
@@ -1,5 +1,6 @@
// 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
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
@@ -31,16 +32,16 @@ void test_div_scale_f32(global float* out, global int* flagout, float a, float b
// 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)
+void test_div_fmas_f32(global float* out, float a, float b, float c, int d)
{
- *out = __builtin_amdgpu_div_fmasf(a, b, c);
+ *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)
+void test_div_fmas_f64(global double* out, double a, double b, double c, int d)
{
- *out = __builtin_amdgpu_div_fmas(a, b, c);
+ *out = __builtin_amdgpu_div_fmas(a, b, c, d);
}
// CHECK-LABEL: @test_div_fixup_f32
@@ -112,3 +113,31 @@ 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);
+}
diff --git a/test/CodeGenOpenCL/const-str-array-decay.cl b/test/CodeGenOpenCL/const-str-array-decay.cl
new file mode 100644
index 000000000000..dbbe08989cb1
--- /dev/null
+++ b/test/CodeGenOpenCL/const-str-array-decay.cl
@@ -0,0 +1,11 @@
+// RUN: %clang_cc1 %s -emit-llvm -o - -ffake-address-space-map | FileCheck %s
+
+int test_func(constant char* foo);
+
+kernel void str_array_decy() {
+ test_func("Test string literal");
+}
+
+// CHECK: i8 addrspace(3)* getelementptr inbounds ([20 x i8] addrspace(3)*
+// CHECK-NOT: addrspacecast
+
diff --git a/test/CodeGenOpenCL/constant-addr-space-globals.cl b/test/CodeGenOpenCL/constant-addr-space-globals.cl
new file mode 100644
index 000000000000..92fb9790b5fb
--- /dev/null
+++ b/test/CodeGenOpenCL/constant-addr-space-globals.cl
@@ -0,0 +1,8 @@
+// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s
+
+// CHECK: @array = addrspace({{[0-9]+}}) constant
+__constant float array[2] = {0.0f, 1.0f};
+
+kernel void test(global float *out) {
+ *out = array[0];
+}
diff --git a/test/CodeGenOpenCL/denorms-are-zero.cl b/test/CodeGenOpenCL/denorms-are-zero.cl
new file mode 100644
index 000000000000..488004fa4fc3
--- /dev/null
+++ b/test/CodeGenOpenCL/denorms-are-zero.cl
@@ -0,0 +1,5 @@
+// RUN: %clang_cc1 -S -cl-denorms-are-zero -o - %s 2>&1
+
+// This test just checks that the -cl-denorms-are-zero argument is accepted
+// by clang. This option is currently a no-op, which is allowed by the
+// OpenCL specification.
diff --git a/test/CodeGenOpenCL/fpmath.cl b/test/CodeGenOpenCL/fpmath.cl
index 704fcd739ac4..ef4da845529c 100644
--- a/test/CodeGenOpenCL/fpmath.cl
+++ b/test/CodeGenOpenCL/fpmath.cl
@@ -22,4 +22,4 @@ double dpscalardiv(double a, double b) {
return a / b;
}
-// CHECK: ![[MD]] = metadata !{float 2.500000e+00}
+// CHECK: ![[MD]] = !{float 2.500000e+00}
diff --git a/test/CodeGenOpenCL/kernel-arg-info.cl b/test/CodeGenOpenCL/kernel-arg-info.cl
index 9832604b0459..4bc191e1d75f 100644
--- a/test/CodeGenOpenCL/kernel-arg-info.cl
+++ b/test/CodeGenOpenCL/kernel-arg-info.cl
@@ -1,28 +1,55 @@
-// RUN: %clang_cc1 %s -cl-kernel-arg-info -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s
+// 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
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_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"}
+// 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"}
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 1, i32 1, i32 1}
-// 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"}
+// 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"}
kernel void foo3(__global half * X) {
}
-// CHECK: metadata !{metadata !"kernel_arg_addr_space", i32 1}
-// CHECK: metadata !{metadata !"kernel_arg_access_qual", metadata !"none"}
-// CHECK: metadata !{metadata !"kernel_arg_type", metadata !"half*"}
-// CHECK: metadata !{metadata !"kernel_arg_type_qual", metadata !""}
-// CHECK: metadata !{metadata !"kernel_arg_name", metadata !"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"}
+
+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"}
+
+typedef image1d_t myImage;
+kernel void foo5(read_only 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"}
diff --git a/test/CodeGenOpenCL/kernel-attributes.cl b/test/CodeGenOpenCL/kernel-attributes.cl
index 0825ffc0336f..8f22d611b820 100644
--- a/test/CodeGenOpenCL/kernel-attributes.cl
+++ b/test/CodeGenOpenCL/kernel-attributes.cl
@@ -8,9 +8,9 @@ kernel __attribute__((vec_type_hint(uint4))) __attribute__((work_group_size_hint
// CHECK: opencl.kernels = !{[[MDNODE0:![0-9]+]], [[MDNODE3:![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 [[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}
+// 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}
diff --git a/test/CodeGenOpenCL/kernel-metadata.cl b/test/CodeGenOpenCL/kernel-metadata.cl
index 3e10a119d044..ef3758fccaa0 100644
--- a/test/CodeGenOpenCL/kernel-metadata.cl
+++ b/test/CodeGenOpenCL/kernel-metadata.cl
@@ -7,4 +7,9 @@ __kernel void kernel_function() {
}
// CHECK: !opencl.kernels = !{!0}
-// CHECK: !0 = metadata !{void ()* @kernel_function}
+// 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"}
diff --git a/test/CodeGenOpenCL/local-initializer-undef.cl b/test/CodeGenOpenCL/local-initializer-undef.cl
new file mode 100644
index 000000000000..5d34f56b821d
--- /dev/null
+++ b/test/CodeGenOpenCL/local-initializer-undef.cl
@@ -0,0 +1,24 @@
+// RUN: %clang_cc1 %s -O0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s
+
+typedef struct Foo {
+ int x;
+ float y;
+ float z;
+} Foo;
+
+// CHECK-DAG: @test.lds_int = internal addrspace(2) global i32 undef
+// CHECK-DAG: @test.lds_int_arr = internal addrspace(2) global [128 x i32] undef
+// CHECK-DAG: @test.lds_struct = internal addrspace(2) global %struct.Foo undef
+// CHECK-DAG: @test.lds_struct_arr = internal addrspace(2) global [64 x %struct.Foo] undef
+__kernel void test()
+{
+ __local int lds_int;
+ __local int lds_int_arr[128];
+ __local Foo lds_struct;
+ __local Foo lds_struct_arr[64];
+
+ lds_int = 1;
+ lds_int_arr[0] = 1;
+ lds_struct.x = 1;
+ lds_struct_arr[0].x = 1;
+}
diff --git a/test/CodeGenOpenCL/local.cl b/test/CodeGenOpenCL/local.cl
index 895c8fa2717d..f1031cdfd84e 100644
--- a/test/CodeGenOpenCL/local.cl
+++ b/test/CodeGenOpenCL/local.cl
@@ -1,9 +1,11 @@
// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - | FileCheck %s
+void func(local int*);
+
__kernel void foo(void) {
- // CHECK: @foo.i = internal unnamed_addr addrspace(2)
+ // CHECK: @foo.i = internal addrspace(2) global i32 undef
__local int i;
- ++i;
+ func(&i);
}
// CHECK-LABEL: define void @_Z3barPU7CLlocali
diff --git a/test/CodeGenOpenCL/opencl_types.cl b/test/CodeGenOpenCL/opencl_types.cl
index 7e99fc548ec5..71d1189e5f08 100644
--- a/test/CodeGenOpenCL/opencl_types.cl
+++ b/test/CodeGenOpenCL/opencl_types.cl
@@ -1,7 +1,7 @@
// RUN: %clang_cc1 %s -emit-llvm -o - -O0 | FileCheck %s
constant sampler_t glb_smp = 7;
-// CHECK: global i32 7
+// CHECK: constant i32 7
void fnc1(image1d_t img) {}
// CHECK: @fnc1(%opencl.image1d_t*
@@ -35,3 +35,6 @@ kernel void foo(image1d_t img) {
fnc4smp(glb_smp);
// 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@@YAXPE?APAUocl_image1d@@PE?APAUocl_image2d@@1@Z"}}
diff --git a/test/CodeGenOpenCL/ptx-calls.cl b/test/CodeGenOpenCL/ptx-calls.cl
index 00f2a0e4c042..bde00bc3d73a 100644
--- a/test/CodeGenOpenCL/ptx-calls.cl
+++ b/test/CodeGenOpenCL/ptx-calls.cl
@@ -9,5 +9,5 @@ __kernel void kernel_function() {
}
// CHECK-LABEL: define void @kernel_function()
// CHECK: call void @device_function()
-// CHECK: !{{[0-9]+}} = metadata !{void ()* @kernel_function, metadata !"kernel", i32 1}
+// CHECK: !{{[0-9]+}} = !{void ()* @kernel_function, !"kernel", i32 1}
diff --git a/test/CodeGenOpenCL/ptx-kernels.cl b/test/CodeGenOpenCL/ptx-kernels.cl
index 49d207f023c8..fc6de4f3d517 100644
--- a/test/CodeGenOpenCL/ptx-kernels.cl
+++ b/test/CodeGenOpenCL/ptx-kernels.cl
@@ -8,4 +8,4 @@ __kernel void kernel_function() {
}
// CHECK-LABEL: define void @kernel_function()
-// CHECK: !{{[0-9]+}} = metadata !{void ()* @kernel_function, metadata !"kernel", i32 1}
+// CHECK: !{{[0-9]+}} = !{void ()* @kernel_function, !"kernel", i32 1}
diff --git a/test/CodeGenOpenCL/relaxed-fpmath.cl b/test/CodeGenOpenCL/relaxed-fpmath.cl
new file mode 100644
index 000000000000..4222ea9c74b5
--- /dev/null
+++ b/test/CodeGenOpenCL/relaxed-fpmath.cl
@@ -0,0 +1,36 @@
+// RUN: %clang_cc1 %s -emit-llvm -o - | FileCheck %s -check-prefix=NORMAL
+// RUN: %clang_cc1 %s -emit-llvm -cl-fast-relaxed-math -o - | FileCheck %s -check-prefix=FAST
+// RUN: %clang_cc1 %s -emit-llvm -cl-finite-math-only -o - | FileCheck %s -check-prefix=FINITE
+// RUN: %clang_cc1 %s -emit-llvm -cl-unsafe-math-optimizations -o - | FileCheck %s -check-prefix=UNSAFE
+// RUN: %clang_cc1 %s -emit-llvm -cl-no-signed-zeros -o - | FileCheck %s -check-prefix=NOSZ
+
+typedef __attribute__(( ext_vector_type(4) )) float float4;
+
+float spscalardiv(float a, float b) {
+ // CHECK: @spscalardiv(
+
+ // NORMAL: fdiv float
+ // FAST: fdiv fast float
+ // FINITE: fdiv nnan ninf float
+ // UNSAFE: fdiv nnan float
+ // NOSZ: fdiv nsz float
+ return a / b;
+}
+// CHECK: attributes
+
+// NORMAL: "no-infs-fp-math"="false"
+// NORMAL: "no-nans-fp-math"="false"
+// NORMAL: "unsafe-fp-math"="false"
+
+// FAST: "no-infs-fp-math"="true"
+// FAST: "no-nans-fp-math"="true"
+// FAST: "unsafe-fp-math"="true"
+
+// FINITE: "no-infs-fp-math"="true"
+// FINITE: "no-nans-fp-math"="true"
+// FINITE: "unsafe-fp-math"="false"
+
+// UNSAFE: "no-infs-fp-math"="false"
+// UNSAFE: "no-nans-fp-math"="true"
+// UNSAFE: "unsafe-fp-math"="true"
+
diff --git a/test/CodeGenOpenCL/str_literals.cl b/test/CodeGenOpenCL/str_literals.cl
index 43c90f83f6e7..092b6372a414 100644
--- a/test/CodeGenOpenCL/str_literals.cl
+++ b/test/CodeGenOpenCL/str_literals.cl
@@ -5,5 +5,5 @@ __constant char * __constant y = "hello world";
// CHECK: unnamed_addr addrspace(3) constant
// CHECK-NOT: addrspace(3) unnamed_addr constant
-// CHECK: @x = addrspace(3) global i8 addrspace(3)*
-// CHECK: @y = addrspace(3) global i8 addrspace(3)*
+// CHECK: @x = addrspace(3) constant i8 addrspace(3)*
+// CHECK: @y = addrspace(3) constant i8 addrspace(3)*
generated by cgit v1.2.3 (git 2.25.1) at 2024-10-30 23:34:03 +0000