Invalid SPIR-V generated from local memory pointer bitcasts

[jrprice]

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]

This is a problem with replace pointer bitcasts. The load is ok, but the store requires an atomic xor to implement correctly.

[andreperezmaselco]

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]

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

[andreperezmaselco]

I'll file this issue, @alan-baker. Thank you.

[rjodinchr]

This is now passing