google / google/clspv

Invalid SPIR-V generated from selection of __constant pointers

Open
#211 0 comments 0 reactions 0 assignees View on GitHub
Dominant language
LLVM
Stars
729
Forks
110
Avg merge
17h 51m
Merged PRs (30d)
23

Description

```
__constant int const_array[] = {0,1,2,3};

void bar(__global int* data, __constant int* c1, __constant int* c2, int x) {
__constant int* c = (x == 0) ? c1 : c2;
*data = *c;
}

__kernel void foo(__global int* data, __constant int* const_arg, int x) {
bar(data, const_array, const_arg, x);
}
```

clspv generates `const_array` as a variable in the `Private` storage class. The ternary statement in `bar` attempts to produce a `StorageBuffer` storage class pointer by selecting between a `Private` storage class pointer and `StorageBuffer` storage class pointer. This is invalid SPIR-V.

The resolution of this issue should also work with #208.

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.