summaryrefslogtreecommitdiffstats
path: root/clang/test/CodeGenOpenCL
diff options
context:
space:
mode:
authorJF Bastien <jfbastien@apple.com>2018-11-15 00:19:18 +0000
committerJF Bastien <jfbastien@apple.com>2018-11-15 00:19:18 +0000
commit3a881e6bbcb06a7e753f88249a29d427ecce64c5 (patch)
tree0017ed6e8adc60a88f22a6f5ef9a71d1ceaeac79 /clang/test/CodeGenOpenCL
parent22bdb3310809d8d2055afc262a0aa81c06100447 (diff)
downloadbcm5719-llvm-3a881e6bbcb06a7e753f88249a29d427ecce64c5.tar.gz
bcm5719-llvm-3a881e6bbcb06a7e753f88249a29d427ecce64c5.zip
CGDecl::emitStoresForConstant fix synthesized constant's name
Summary: The name of the synthesized constants for constant initialization was using mangling for statics, which isn't generally correct and (in a yet-uncommitted patch) causes the mangler to assert out because the static ends up trying to mangle function parameters and this makes no sense. Instead, mangle to `"__const." + FunctionName + "." + DeclName`. Reviewers: rjmccall Subscribers: dexonsmith, cfe-commits Differential Revision: https://reviews.llvm.org/D54055 llvm-svn: 346915
Diffstat (limited to 'clang/test/CodeGenOpenCL')
-rw-r--r--clang/test/CodeGenOpenCL/amdgpu-nullptr.cl4
-rw-r--r--clang/test/CodeGenOpenCL/constant-addr-space-globals.cl4
-rw-r--r--clang/test/CodeGenOpenCL/partial_initializer.cl4
-rw-r--r--clang/test/CodeGenOpenCL/private-array-initialization.cl4
4 files changed, 8 insertions, 8 deletions
diff --git a/clang/test/CodeGenOpenCL/amdgpu-nullptr.cl b/clang/test/CodeGenOpenCL/amdgpu-nullptr.cl
index 688d3a58e90..c7c77920b77 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-nullptr.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-nullptr.cl
@@ -143,7 +143,7 @@ void test_static_var_local(void) {
// NOOPT: store i8 addrspace(5)* null, i8 addrspace(5)* addrspace(5)* %sp3, align 4
// NOOPT: store i8 addrspace(5)* null, i8 addrspace(5)* addrspace(5)* %sp4, align 4
// NOOPT: %[[SS1:.*]] = bitcast %struct.StructTy1 addrspace(5)* %SS1 to i8 addrspace(5)*
-// NOOPT: call void @llvm.memcpy.p5i8.p4i8.i64(i8 addrspace(5)* align 8 %[[SS1]], i8 addrspace(4)* align 8 bitcast (%struct.StructTy1 addrspace(4)* @test_func_scope_var_private.SS1 to i8 addrspace(4)*), i64 32, i1 false)
+// NOOPT: call void @llvm.memcpy.p5i8.p4i8.i64(i8 addrspace(5)* align 8 %[[SS1]], i8 addrspace(4)* align 8 bitcast (%struct.StructTy1 addrspace(4)* @__const.test_func_scope_var_private.SS1 to i8 addrspace(4)*), i64 32, i1 false)
// NOOPT: %[[SS2:.*]] = bitcast %struct.StructTy2 addrspace(5)* %SS2 to i8 addrspace(5)*
// NOOPT: call void @llvm.memset.p5i8.i64(i8 addrspace(5)* align 8 %[[SS2]], i8 0, i64 24, i1 false)
void test_func_scope_var_private(void) {
@@ -163,7 +163,7 @@ void test_func_scope_var_private(void) {
// NOOPT: store i8 addrspace(3)* null, i8 addrspace(3)* addrspace(5)* %sp3, align 4
// NOOPT: store i8 addrspace(3)* null, i8 addrspace(3)* addrspace(5)* %sp4, align 4
// NOOPT: %[[SS1:.*]] = bitcast %struct.StructTy1 addrspace(5)* %SS1 to i8 addrspace(5)*
-// NOOPT: call void @llvm.memcpy.p5i8.p4i8.i64(i8 addrspace(5)* align 8 %[[SS1]], i8 addrspace(4)* align 8 bitcast (%struct.StructTy1 addrspace(4)* @test_func_scope_var_local.SS1 to i8 addrspace(4)*), i64 32, i1 false)
+// NOOPT: call void @llvm.memcpy.p5i8.p4i8.i64(i8 addrspace(5)* align 8 %[[SS1]], i8 addrspace(4)* align 8 bitcast (%struct.StructTy1 addrspace(4)* @__const.test_func_scope_var_local.SS1 to i8 addrspace(4)*), i64 32, i1 false)
// NOOPT: %[[SS2:.*]] = bitcast %struct.StructTy2 addrspace(5)* %SS2 to i8 addrspace(5)*
// NOOPT: call void @llvm.memset.p5i8.i64(i8 addrspace(5)* align 8 %[[SS2]], i8 0, i64 24, i1 false)
void test_func_scope_var_local(void) {
diff --git a/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl b/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl
index 7bb970527c2..5fcf117dde3 100644
--- a/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl
+++ b/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl
@@ -13,8 +13,8 @@ kernel void test(global float *out) {
void foo(constant int* p, constant const int *p1, const int *p2, const int *p3);
// CHECK: @k.arr1 = internal addrspace(2) constant [3 x i32] [i32 1, i32 2, i32 3]
-// CHECK: @k.arr2 = private unnamed_addr addrspace(2) constant [3 x i32] [i32 4, i32 5, i32 6]
-// CHECK: @k.arr3 = private unnamed_addr addrspace(2) constant [3 x i32] [i32 7, i32 8, i32 9]
+// CHECK: @__const.k.arr2 = private unnamed_addr addrspace(2) constant [3 x i32] [i32 4, i32 5, i32 6]
+// CHECK: @__const.k.arr3 = private unnamed_addr addrspace(2) constant [3 x i32] [i32 7, i32 8, i32 9]
// CHECK: @k.var1 = internal addrspace(2) constant i32 1
kernel void k(void) {
// CHECK-NOT: %arr1 = alloca [3 x i32]
diff --git a/clang/test/CodeGenOpenCL/partial_initializer.cl b/clang/test/CodeGenOpenCL/partial_initializer.cl
index ee6be919a7f..4e829958413 100644
--- a/clang/test/CodeGenOpenCL/partial_initializer.cl
+++ b/clang/test/CodeGenOpenCL/partial_initializer.cl
@@ -24,7 +24,7 @@ int4 GV1 = (int4)((int2)(1,2),3,4);
// CHECK: @GV2 = addrspace(1) global <4 x i32> <i32 1, i32 1, i32 1, i32 1>, align 16
int4 GV2 = (int4)(1);
-// CHECK: @f.S = private unnamed_addr addrspace(2) constant %struct.StrucTy { i32 1, i32 2, i32 0 }, align 4
+// CHECK: @__const.f.S = private unnamed_addr addrspace(2) constant %struct.StrucTy { i32 1, i32 2, i32 0 }, align 4
// CHECK-LABEL: define spir_func void @f()
void f(void) {
@@ -46,7 +46,7 @@ void f(void) {
float A[6][6] = {1.0f, 2.0f};
// CHECK: %[[v5:.*]] = bitcast %struct.StrucTy* %S to i8*
- // CHECK: call void @llvm.memcpy.p0i8.p2i8.i32(i8* align 4 %[[v5]], i8 addrspace(2)* align 4 bitcast (%struct.StrucTy addrspace(2)* @f.S to i8 addrspace(2)*), i32 12, i1 false)
+ // CHECK: call void @llvm.memcpy.p0i8.p2i8.i32(i8* align 4 %[[v5]], i8 addrspace(2)* align 4 bitcast (%struct.StrucTy addrspace(2)* @__const.f.S to i8 addrspace(2)*), i32 12, i1 false)
StrucTy S = {1, 2};
// CHECK: store <2 x i32> <i32 1, i32 2>, <2 x i32>* %[[compoundliteral1]], align 8
diff --git a/clang/test/CodeGenOpenCL/private-array-initialization.cl b/clang/test/CodeGenOpenCL/private-array-initialization.cl
index 9aa058dcfac..420270de193 100644
--- a/clang/test/CodeGenOpenCL/private-array-initialization.cl
+++ b/clang/test/CodeGenOpenCL/private-array-initialization.cl
@@ -6,11 +6,11 @@
void test() {
__private int arr[] = {1, 2, 3};
// PRIVATE0: %[[arr_i8_ptr:[0-9]+]] = bitcast [3 x i32]* %arr to i8*
-// PRIVATE0: call void @llvm.memcpy.p0i8.p2i8.i32(i8* align 4 %[[arr_i8_ptr]], i8 addrspace(2)* align 4 bitcast ([3 x i32] addrspace(2)* @test.arr to i8 addrspace(2)*), i32 12, i1 false)
+// PRIVATE0: call void @llvm.memcpy.p0i8.p2i8.i32(i8* align 4 %[[arr_i8_ptr]], i8 addrspace(2)* align 4 bitcast ([3 x i32] addrspace(2)* @__const.test.arr to i8 addrspace(2)*), i32 12, i1 false)
// PRIVATE5: %arr = alloca [3 x i32], align 4, addrspace(5)
// PRIVATE5: %0 = bitcast [3 x i32] addrspace(5)* %arr to i8 addrspace(5)*
-// PRIVATE5: call void @llvm.memcpy.p5i8.p4i8.i64(i8 addrspace(5)* align 4 %0, i8 addrspace(4)* align 4 bitcast ([3 x i32] addrspace(4)* @test.arr to i8 addrspace(4)*), i64 12, i1 false)
+// PRIVATE5: call void @llvm.memcpy.p5i8.p4i8.i64(i8 addrspace(5)* align 4 %0, i8 addrspace(4)* align 4 bitcast ([3 x i32] addrspace(4)* @__const.test.arr to i8 addrspace(4)*), i64 12, i1 false)
}
__kernel void initializer_cast_is_valid_crash() {
OpenPOWER on IntegriCloud