LLVM-SPIRV crashes on struct with vector-content
Opened this issue · 2 comments
doe300 commented
The original OpenCL-C code:
struct SomeStruct
{
float16 f;
int i;
};
__kernel void test_struct(const __global struct SomeStruct* in, __global struct SomeStruct* out)
{
struct SomeStruct tmp = *in;
tmp.i = 42;
*out = tmp;
}
The LLVM-IR code (excerpt), compiled with clang version 3.6.1 (https://github.com/KhronosGroup/SPIR d7e44c3b27581e54ca0e522987d1ade2bd29b70d) (https://github.com/KhronosGroup/SPIRV-LLVM.git fffc52b7fb3552c045c398d68e22216c8d770c00):
%struct.SomeStruct = type { <16 x float>, i32, [60 x i8] }
define spir_kernel void @test_struct(%struct.SomeStruct addrspace(1)* nocapture readonly %in, %struct.SomeStruct addrspace(1)* nocapture %out) #0 {
%tmp.sroa.5 = alloca [60 x i8], align 4
%1 = getelementptr inbounds [60 x i8]* %tmp.sroa.5, i32 0, i32 0
call void @llvm.lifetime.start(i64 60, i8* %1)
%2 = getelementptr inbounds %struct.SomeStruct addrspace(1)* %in, i32 0, i32 0
%3 = load <16 x float> addrspace(1)* %2, align 64
%4 = getelementptr inbounds %struct.SomeStruct addrspace(1)* %in, i32 0, i32 2, i32 0
call void @llvm.memcpy.p0i8.p1i8.i32(i8* %1, i8 addrspace(1)* %4, i32 60, i32 4, i1 false)
%5 = getelementptr inbounds %struct.SomeStruct addrspace(1)* %out, i32 0, i32 0
store <16 x float> %3, <16 x float> addrspace(1)* %5, align 64
%6 = getelementptr inbounds %struct.SomeStruct addrspace(1)* %out, i32 0, i32 1
store i32 42, i32 addrspace(1)* %6, align 64
%7 = getelementptr inbounds %struct.SomeStruct addrspace(1)* %out, i32 0, i32 2, i32 0
call void @llvm.memcpy.p1i8.p0i8.i32(i8 addrspace(1)* %7, i8* %1, i32 60, i32 4, i1 false)
call void @llvm.lifetime.end(i64 60, i8* %1)
ret void
}
declare void @llvm.lifetime.start(i64, i8* nocapture)
declare void @llvm.memcpy.p0i8.p1i8.i32(i8* nocapture, i8 addrspace(1)* nocapture readonly, i32, i32, i1)
declare void @llvm.memcpy.p1i8.p0i8.i32(i8 addrspace(1)* nocapture, i8* nocapture readonly, i32, i32, i1)
declare void @llvm.lifetime.end(i64, i8* nocapture)
llvm-spirv crashes with following stack-trace:
Stack dump:
0. Program arguments: /opt/SPIRV-LLVM/build/bin/llvm-spirv -o /tmp/out.spv
0 llvm-spirv 0x00000000005f5972
1 llvm-spirv 0x00000000005f4241
2 libpthread.so.0 0x00007f09af1ab5c0
3 llvm-spirv 0x00000000004b0d70
4 llvm-spirv 0x00000000004b1683
5 llvm-spirv 0x00000000004acefd
6 llvm-spirv 0x00000000004ad256
7 llvm-spirv 0x00000000004b1f08
8 llvm-spirv 0x00000000004b2240
9 llvm-spirv 0x00000000004b226b
10 llvm-spirv 0x000000000058bcac
11 llvm-spirv 0x00000000004b22ce
12 llvm-spirv 0x0000000000407a73
13 libc.so.6 0x00007f09ae54c401 __libc_start_main + 241
14 llvm-spirv 0x000000000040abea
Segmentation fault (core dumped)
Analysis with valgrind prints this:
Process terminating with default action of signal 11 (SIGSEGV)
Access not within mapped region at address 0x8
at 0x4CED70: SPIRV::SPIRVEntry::setModule(SPIRV::SPIRVModule*) (in /opt/SPIRV-LLVM/build/bin/llvm-spirv)
by 0x45F298: SPIRV::SPIRVDecoder::getEntry() (in /opt/SPIRV-LLVM/build/bin/llvm-spirv)
by 0x44793F: SPIRV::operator>>(std::istream&, SPIRV::SPIRVModule&) (in /opt/SPIRV-LLVM/build/bin/llvm-spirv)
by 0x458F13: SPIRV::ConvertSPIRV(std::istream&, llvm::raw_ostream&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >&, bool, bool) (in /opt/SPIRV-LLVM/build/bin/llvm-spirv)
by 0x40AFBB: convertSPIRV()::{lambda(llvm::raw_ostream&)#1}::operator()(llvm::raw_ostream&) const (in /opt/SPIRV-LLVM/build/bin/llvm-spirv)
by 0x40B2DB: convertSPIRV() (in /opt/SPIRV-LLVM/build/bin/llvm-spirv)
by 0x40798B: main (in /opt/SPIRV-LLVM/build/bin/llvm-spirv)
... which looks like access to a nullptr
AlexeySotkin commented
Could you try with debug build (with enabled assert functions) of llvm-spirv?
This code worked for me, when I commented out calls to llvm.lifetime.* intrinsics.
doe300 commented
You are right, seems like the llvm.lifetime.* intrinsics arethe culprits.
Here the new valgrind output:
Process terminating with default action of signal 6 (SIGABRT)
at 0x5D4191F: raise (raise.c:58)
by 0x5D43632: abort (abort.c:118)
by 0x5D39DA6: __assert_fail_base (assert.c:92)
by 0x5D39E51: __assert_fail (assert.c:101)
by 0x56AA54: SPIRV::LLVMToSPIRV::transLifetimeIntrinsicInst(spv::Op, llvm::IntrinsicInst*, SPIRV::SPIRVBasicBlock*) (SPIRVWriter.cpp:855)
by 0x56DF1D: SPIRV::LLVMToSPIRV::transIntrinsicInst(llvm::IntrinsicInst*, SPIRV::SPIRVBasicBlock*) (SPIRVWriter.cpp:1307)
by 0x56D17B: SPIRV::LLVMToSPIRV::transValueWithoutDecoration(llvm::Value*, SPIRV::SPIRVBasicBlock*, bool) (SPIRVWriter.cpp:1157)
by 0x56A5DF: SPIRV::LLVMToSPIRV::transValue(llvm::Value*, SPIRV::SPIRVBasicBlock*, bool) (SPIRVWriter.cpp:812)
by 0x56ED83: SPIRV::LLVMToSPIRV::transFunction(llvm::Function*) (SPIRVWriter.cpp:1464)
by 0x56F2E3: SPIRV::LLVMToSPIRV::translate() (SPIRVWriter.cpp:1508)
by 0x5663E3: SPIRV::LLVMToSPIRV::runOnModule(llvm::Module&) (SPIRVWriter.cpp:178)
by 0x6CF415: (anonymous namespace)::MPPassManager::runOnModule(llvm::Module&) (LegacyPassManager.cpp:1616)