Segfault when compiling constant structs with pointers to constant structs
I'm working on some unnamed spir-v tooling, and hit a really tough case to compile. So, I figured I'd look at other tools and see how they handle it... and now I'm submitting this! :P
__constant static struct A {
int f;
} a = { 1 };
__constant static struct B {
__constant struct A* f;
} b = { &a };
__constant static struct C {
__constant struct B* f;
} c = { &b };
__kernel void foo(__global int* data) {
*data += c.f->f->f;
}
when compiled, produces a segfault.
> clspv asdf.cl -o asdf.spv
<badref> = internal addrspace(2) constant %struct.A fish: “clspv asdf.cl -o asdf.spv” terminated by signal SIGSEGV (Address boundary error)
Something to note is that you need three structs to trigger the error - if you change the entry point to *data += b.f->f;, no segfault occurs (and a "c is unused" warning).
For some additional context, SPIRV-LLVM-Translator also fails to compile the above sample (with clang -cc1 -triple spir -cl-std=CL1.2 asdf.cl -O0 -emit-llvm-bc -o test.bc && llvm-spirv test.bc -o test.spv) - but it fails by producing invalid spir-v, it includes an OpVariable reference in an OpConstantComposite.
Thanks for the bug report. Unfortunately, clspv can't handle this case due to restrictions on shaders. This particular case might be workable, but it quickly becomes intractable. That said, the compiler should produce a real error message.
This is not segfaulting anymore and produces correct spirv.