summaryrefslogtreecommitdiffstats
path: root/test/CodeGenOpenCL
diff options
context:
space:
mode:
authorRafael Espindola <rafael.espindola@gmail.com>2018-02-23 19:30:48 +0000
committerRafael Espindola <rafael.espindola@gmail.com>2018-02-23 19:30:48 +0000
commit0e355701be41ad5bc6f83313a7f1ced24336c42d (patch)
treee6d410c6da69daa13c8ad94d2b32259d58fb976f /test/CodeGenOpenCL
parentc036db5904d43222dd47aae0b1a3e5a2e0c44ad6 (diff)
Bring r325915 back.
The tests that failed on a windows host have been fixed. Original message: Start setting dso_local for COFF. With this there are still some GVs where we don't set dso_local because setGVProperties is never called. I intend to fix that in followup commits. This is just the bare minimum to teach shouldAssumeDSOLocal what it should do for COFF. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@325940 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGenOpenCL')
-rw-r--r--test/CodeGenOpenCL/address-spaces.cl18
-rw-r--r--test/CodeGenOpenCL/kernel-attributes.cl6
-rw-r--r--test/CodeGenOpenCL/kernel-metadata.cl2
-rw-r--r--test/CodeGenOpenCL/str_literals.cl4
4 files changed, 15 insertions, 15 deletions
diff --git a/test/CodeGenOpenCL/address-spaces.cl b/test/CodeGenOpenCL/address-spaces.cl
index f4ddbc611c..60f5e30d32 100644
--- a/test/CodeGenOpenCL/address-spaces.cl
+++ b/test/CodeGenOpenCL/address-spaces.cl
@@ -13,13 +13,13 @@ struct S {
int *z;
};
-// CL20-DAG: @g_extern_var = external addrspace(1) global float
-// CL20-DAG: @l_extern_var = external addrspace(1) global float
+// CL20-DAG: @g_extern_var = external {{(dso_local )?}}addrspace(1) global float
+// CL20-DAG: @l_extern_var = external {{(dso_local )?}}addrspace(1) global float
// CL20-DAG: @test_static.l_static_var = internal addrspace(1) global float 0.000000e+00
// CL20-DAG: @g_static_var = internal addrspace(1) global float 0.000000e+00
#ifdef CL20
-// CL20-DAG: @g_s = common addrspace(1) global %struct.S zeroinitializer
+// CL20-DAG: @g_s = common {{(dso_local )?}}addrspace(1) global %struct.S zeroinitializer
struct S g_s;
#endif
@@ -53,10 +53,10 @@ void fc(constant int *arg) {}
#ifdef CL20
int i;
-// CL20-DAG: @i = common addrspace(1) global i32 0
+// CL20-DAG: @i = common {{(dso_local )?}}addrspace(1) global i32 0
int *ptr;
-// CL20SPIR-DAG: @ptr = common addrspace(1) global i32 addrspace(4)* null
-// CL20AMDGCN-DAG: @ptr = common addrspace(1) global i32* null
+// CL20SPIR-DAG: @ptr = common {{(dso_local )?}}addrspace(1) global i32 addrspace(4)* null
+// CL20AMDGCN-DAG: @ptr = common {{(dso_local )?}}addrspace(1) global i32* null
#endif
// SPIR: i32* %arg
@@ -79,13 +79,13 @@ void f(int *arg) {
typedef int int_td;
typedef int *intp_td;
-// SPIR: define void @test_typedef(i32 addrspace(1)* %x, i32 addrspace(2)* %y, i32* %z)
+// SPIR: define {{(dso_local )?}}void @test_typedef(i32 addrspace(1)* %x, i32 addrspace(2)* %y, i32* %z)
void test_typedef(global int_td *x, constant int_td *y, intp_td z) {
*x = *y;
*z = 0;
}
-// SPIR: define void @test_struct()
+// SPIR: define {{(dso_local )?}}void @test_struct()
void test_struct() {
// SPIR: %ps = alloca %struct.S*
// CL20SPIR: %ps = alloca %struct.S addrspace(4)*
@@ -99,7 +99,7 @@ void test_struct() {
#endif
}
-// SPIR-LABEL: define void @test_void_par()
+// SPIR-LABEL: define {{(dso_local )?}}void @test_void_par()
void test_void_par(void) {}
// On ppc64 returns signext i32.
diff --git a/test/CodeGenOpenCL/kernel-attributes.cl b/test/CodeGenOpenCL/kernel-attributes.cl
index eff2e57d8c..c9ecb144c9 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 spir_kernel void @kernel1(i32 {{[^%]*}}%a) {{[^{]+}} !vec_type_hint ![[MD1:[0-9]+]] !reqd_work_group_size ![[MD2:[0-9]+]]
+// CHECK: define {{(dso_local )?}}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 spir_kernel void @kernel2(i32 {{[^%]*}}%a) {{[^{]+}} !vec_type_hint ![[MD3:[0-9]+]] !work_group_size_hint ![[MD4:[0-9]+]]
+// CHECK: define {{(dso_local )?}}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 spir_kernel void @kernel3(i32 {{[^%]*}}%a) {{[^{]+}} !intel_reqd_sub_group_size ![[MD5:[0-9]+]]
+// CHECK: define {{(dso_local )?}}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 95be43015a..cdec97b7df 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 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: 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/str_literals.cl b/test/CodeGenOpenCL/str_literals.cl
index 8fe5cefc8d..436af83afe 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(2) constant
// CHECK-NOT: addrspace(2) unnamed_addr constant
-// CHECK: @x = addrspace(2) constant i8 addrspace(2)*
-// CHECK: @y = addrspace(2) constant i8 addrspace(2)*
+// CHECK: @x = {{(dso_local )?}}addrspace(2) constant i8 addrspace(2)*
+// CHECK: @y = {{(dso_local )?}}addrspace(2) constant i8 addrspace(2)*