KhronosGroup / KhronosGroup/SPIRV-LLVM-Translator
Incorrect spir-v generated when compiling constant structs with pointers to constant structs
- Dominant language
- LLVM
- Stars
- 625
- Forks
- 279
- Avg merge
- 3d 5h
- Merged PRs (30d)
- 34
Description
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
Command: `clang -cc1 -triple spir -cl-std=CL1.2 asdf.cl -O0 -emit-llvm-bc -o test.bc && llvm-spirv test.bc -o test.spv`
```opencl
__constant static struct A {
int f;
} a = { 1 };
__constant static struct B {
__constant struct A* f;
} b = { &a };
__kernel void foo(__global int* data) {
*data += b.f->f;
}
```
when compiled, produces invalid IR.
```
> spirv-val test.spv
error: line 38: OpConstantComposite Constituent '7[%a]' is not a constant or undef.
%9 = OpConstantComposite %struct_B %a
%a = OpVariable %_ptr_UniformConstant_struct_A UniformConstant %5
```
(i.e. referencing an OpVariable from an OpConstantComposite)
Something to note is that if you compile with clang -O2, then clang completely optimizes out the pointer and references the constant `1` directly, and so the tough case doesn't even reach llvm-spirv - so you must compile this particular repro with -O0.
---
For some additional context, [clspv also fails to compile the above sample](https://github.com/google/clspv/issues/631) - but it fails by segfaulting, not by producing invalid spir-v.
Contributor guide
Research direction
Start by running the provided clang and llvm-spirv command on the OpenCL sample, then validate test.spv with spirv-val. The fix is complete when this -O0 case produces valid SPIR-V without an OpConstantComposite referencing the OpVariable %a.
Written by the indexing model from the issue text.
Assessment
- Domain
- compilers
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Stale
- Clarity
- Mostly clear
- Newbie friendliness
- 45/100