KhronosGroup / KhronosGroup/SPIRV-LLVM-Translator

Incorrect spir-v generated when compiling constant structs with pointers to constant structs

Open
#729 0 comments 0 reactions 0 assignees View on GitHub
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.