KhronosGroup / KhronosGroup/SPIRV-LLVM-Translator
Incorrect Op*AccessChain instruction for vec ptr to scalar ptr casts
- Dominant language
- LLVM
- Stars
- 625
- Forks
- 279
- Avg merge
- 3d 5h
- Merged PRs (30d)
- 34
Description
I was debugging an issue w/ vstore_local cl cts tests, which does something like:
```
vstore2( srcValues[ tid ], offsets[ tid ], ( (__local char *)sSharedStorage ) + alignmentOffset );
```
I've simplified this to a more trivial example:
```
__kernel void test_fn(__global int *res, __global int2 *v)
{
*res = ((__global int *)v)[1];
}
```
This generates an `AccessChain` instruction like:
```
%10 = OpFunctionParameter %_ptr_CrossWorkgroup_v2uint
%15 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %10 %uint_0 %uint_1
```
I believe this should have been an `OpInBoundsAccessChain` instead, or equivalently:
```
%15 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %10 %uint_0 %uint_0 %uint_1
```
given that for the Ptr access chain variants:
> The type of Base after being dereferenced with Element is still the same as the original type of Base.
Note also spirv-val doesn't seem to be checking the resulting type of the AccessChain instruction, since it doesn't seem to complain about this. Perhaps a separate bug?
(Originally filed at https://github.com/KhronosGroup/SPIRV-Tools/issues/1498 but that seems to be the wrong place)
Contributor guide
Research direction
Start with the simplified OpenCL reproducer and inspect the generated SPIR-V AccessChain instruction shown in the issue. Compare the emitted pointer access-chain form with the stated expected forms, and check how spirv-val handles the result. Done means the cast produces a valid, type-correct access chain for this case and the behavior is covered by an appropriate regression check.
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
- 35/100