The Wayback Machine - https://web.archive.org/web/20201128060415/https://github.com/google/clspv/issues/532
Skip to content

/ clspv

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Invalid SPIR-V generated from local memory pointer bitcasts #532

Open
jrprice opened this issue Mar 18, 2020 · 4 comments
Open

Invalid SPIR-V generated from local memory pointer bitcasts #532

jrprice opened this issue Mar 18, 2020 · 4 comments
Projects
General

Comments

@jrprice
Copy link
Collaborator

@jrprice jrprice commented Mar 18, 2020

This kernel:

__kernel void foo(
 constant const uchar *indices,
 local int4* __shared)
{
  int x = get_local_id(0);
  short _827 = ((local short *)__shared)[x];
  uchar _833 = indices[_827];
  int y = x + 4872;
  ((local uchar *)__shared)[y] = _833;
}

produces the following SPIR-V validation error:

error: line 75: OpStore Pointer <id> '55[%55]' is not a logical pointer.
  OpStore %55 %50

since the OpStore's pointer operand comes from an OpBitcast instruction.
@alan-baker
Copy link
Collaborator

@alan-baker alan-baker commented Mar 19, 2020

This is a problem with replace pointer bitcasts. The load is ok, but the store requires an atomic xor to implement correctly.
@alan-baker alan-baker added this to Bugs in General Mar 19, 2020
@andreperezmaselco
Copy link

@andreperezmaselco andreperezmaselco commented Mar 22, 2020 •
edited

There is a problem with ReplaceLLVMIntrinsicsPass::replaceMemcpy.
CI->getArgOperand(0) and CI->getArgOperand(1) returns can be a GEPOperator instead of a BitCastOperator. So, some logic must be added to cover the GEPOperator case.
@alan-baker, may I work on this issue to understand more about clspv?
@alan-baker
Copy link
Collaborator

@alan-baker alan-baker commented Mar 23, 2020

@andreperezmaselco that sounds like a separate issue from this. You're welcome to take a look, but please file an issue for it.
@andreperezmaselco
Copy link

@andreperezmaselco andreperezmaselco commented Mar 23, 2020

I'll file this issue, @alan-baker.
Thank you.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
3 participants
@jrprice @andreperezmaselco @alan-baker
You can’t perform that action at this time.