KhronosGroup / KhronosGroup/SPIRV-LLVM-Translator

Incorrect Op*AccessChain instruction for vec ptr to scalar ptr casts

Open
#10 3 comments 0 reactions 0 assignees View on GitHub
bug stale
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

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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.