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

[Zhigang Gong]

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