LLVM-SPIRV crashes on struct with vector-content

[doe300]

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]

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]

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)