diff options
| author | Matt Arsenault <Matthew.Arsenault@amd.com> | 2018-02-08 19:37:09 +0000 |
|---|---|---|
| committer | Matt Arsenault <Matthew.Arsenault@amd.com> | 2018-02-08 19:37:09 +0000 |
| commit | 935574a490be6f258a6b950bbb94b343e87876d2 (patch) | |
| tree | ea7093ee24e7a37c9988a3d81a85d79d77141f2a /clang/test/CodeGenOpenCL/private-array-initialization.cl | |
| parent | 7f7891c427daa063c350f63175e0d186b829d683 (diff) | |
| download | bcm5719-llvm-935574a490be6f258a6b950bbb94b343e87876d2.tar.gz bcm5719-llvm-935574a490be6f258a6b950bbb94b343e87876d2.zip | |
Fix crash on array initializer with non-0 alloca addrspace
llvm-svn: 324641
Diffstat (limited to 'clang/test/CodeGenOpenCL/private-array-initialization.cl')
| -rw-r--r-- | clang/test/CodeGenOpenCL/private-array-initialization.cl | 30 |
1 files changed, 27 insertions, 3 deletions
diff --git a/clang/test/CodeGenOpenCL/private-array-initialization.cl b/clang/test/CodeGenOpenCL/private-array-initialization.cl index 6c2f45f8378..018128091a2 100644 --- a/clang/test/CodeGenOpenCL/private-array-initialization.cl +++ b/clang/test/CodeGenOpenCL/private-array-initialization.cl @@ -1,9 +1,33 @@ -// RUN: %clang_cc1 %s -triple spir-unknown-unknown -O0 -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple spir-unknown-unknown -O0 -emit-llvm -o - | FileCheck -check-prefix=PRIVATE0 %s +// RUN: %clang_cc1 %s -triple amdgcn-amd-amdhsa-unknown -O0 -emit-llvm -o - | FileCheck -check-prefix=PRIVATE5 %s // CHECK: @test.arr = private unnamed_addr addrspace(2) constant [3 x i32] [i32 1, i32 2, i32 3], align 4 void test() { __private int arr[] = {1, 2, 3}; -// CHECK: %[[arr_i8_ptr:[0-9]+]] = bitcast [3 x i32]* %arr to i8* -// CHECK: 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: %[[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) + +// 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.p2i8.i64(i8 addrspace(5)* align 4 %0, i8 addrspace(2)* align 4 bitcast ([3 x i32] addrspace(2)* @test.arr to i8 addrspace(2)*), i64 12, i1 false) +} + +__kernel void initializer_cast_is_valid_crash() { +// PRIVATE0: %v512 = alloca [64 x i8], align 1 +// PRIVATE0: %0 = bitcast [64 x i8]* %v512 to i8* +// PRIVATE0: call void @llvm.memset.p0i8.i32(i8* align 1 %0, i8 0, i32 64, i1 false) +// PRIVATE0: %1 = bitcast i8* %0 to [64 x i8]* + + +// PRIVATE5: %v512 = alloca [64 x i8], align 1, addrspace(5) +// PRIVATE5: %0 = bitcast [64 x i8] addrspace(5)* %v512 to i8 addrspace(5)* +// PRIVATE5: call void @llvm.memset.p5i8.i64(i8 addrspace(5)* align 1 %0, i8 0, i64 64, i1 false) +// PRIVATE5: %1 = bitcast i8 addrspace(5)* %0 to [64 x i8] addrspace(5)* + unsigned char v512[64] = { + 0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00, + 0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00, + 0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00, + 0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x02,0x00 + }; } |
