[Beignet] compiler asserts on unsupported instrinsics compiling kernel with a single function

[Edward Ching]

When compiling a simple single function kernel (patch attached, also cut
and pasted below), the compiler asserted on unsupported intrinsics:

ASSERTION FAILED: Unsupported intrinsics
  at file
/root/WORK/test_split/beignet/backend/src/llvm/llvm_gen_backend.cpp,
function void gbe::GenWriter::regAllocateCallInst(llvm::CallInst&), line
1688
Stack dump:
0. Running pass 'Function Pass Manager' on module '/tmp/file0yciwa.ll'.
1. Running pass 'Gen Back-End' on function '@test_split'

It looks like this is triggered by the "shr_mem[pid+offset]=0;" line when 0
is assigned. Is this a bug? It looks like the kernel code is not doing
anything illegal. It compiled ok using Intel OpenCL SDK 2013 (on Linux, for
CPU)

Thanks,
/Ed

test_split.cl:

__kernel void
test_split(
__global int *in,
int stride,
int n_recs,
int n_parts,
__global int *out,
__local  int *shr_mem)
{
int pid;
int glb_tid = get_global_id(0);
int loc_tid = get_local_id(0);
int offset  = loc_tid*n_parts;

for(pid=0; pid<n_parts; pid++)
{
shr_mem[pid+offset]=0;
}

for(int pos=glb_tid; pos<n_recs; pos+=stride)
{
pid=in[pos];
shr_mem[pid+offset]++;
}

for(pid=0; pid<n_parts; pid++)
{
out[glb_tid + stride*pid] = shr_mem[pid+offset];
}
}

compiler_test_split.cpp:

void compiler_test_split(void)
{
  OCL_CREATE_KERNEL("test_split");
}
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.freedesktop.org/archives/beignet/attachments/20130720/6d48c373/attachment.html>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: 0001-test-case-to-trigger-Unsupported-intrinsics-when-com.patch
Type: application/octet-stream
Size: 2014 bytes
Desc: not available
URL: <http://lists.freedesktop.org/archives/beignet/attachments/20130720/6d48c373/attachment.obj>