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

[Edward Ching]

It compiles fine now with the latest master branch with the tip at
commit 54d31547b9e4ce10e58a6e116ee33b5aba7b4ab8
Author: Yi Sun <yi.sun at intel.com>
Date:   Mon Jul 22 15:59:42 2013 +0800
utest: add built-in test case for get_global_id.

Thanks!
/Ed



On Mon, Jul 22, 2013 at 12:29 AM, Zhigang Gong <zhigang.gong at linux.intel.com
> wrote:

> Hi Edward,
>
> Could you try to compile your kernel again with latest master branch?
> Rong's patch should get this issue fixed.
>
> On Sat, Jul 20, 2013 at 10:49:08PM -0700, Edward Ching wrote:
> > 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");
> > }
>
>
> > _______________________________________________
> > Beignet mailing list
> > Beignet at lists.freedesktop.org
> > http://lists.freedesktop.org/mailman/listinfo/beignet
>
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.freedesktop.org/archives/beignet/attachments/20130722/4d40b3a4/attachment.html>