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

[Zhigang Gong]

Hi Edward,

Thanks for reporting this bug, I did a quick look at the LLVM IR layer of
this kernel. And I think the root cause is that the shr_mem[pid+offset] = 0
in a for loop make the compiler generate a llvm.memset intrinsic which is
not implemented.

We need to implement these intrinsics. Anybody from this list want to take
this work item?

Edward, before we get this intrinsic implemented, if you need a workaround,
here is a simple one, define a constant zero

constant int zero = 0;
...
shr_mem[pid+offset] = zero;

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