diff options
author | JF Bastien <jfbastien@apple.com> | 2018-11-15 00:19:18 +0000 |
---|---|---|
committer | JF Bastien <jfbastien@apple.com> | 2018-11-15 00:19:18 +0000 |
commit | 3a881e6bbcb06a7e753f88249a29d427ecce64c5 (patch) | |
tree | 0017ed6e8adc60a88f22a6f5ef9a71d1ceaeac79 /clang/test/CodeGenOpenCL | |
parent | 22bdb3310809d8d2055afc262a0aa81c06100447 (diff) | |
download | bcm5719-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')
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() { |