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

[Yang, Rong R]

Can we simply add clang build option -fno-builtin to disable these intrinsic?

-----Original Message-----
From: beignet-bounces+rong.r.yang=intel.com at lists.freedesktop.org [mailto:beignet-bounces+rong.r.yang=intel.com at lists.freedesktop.org] On Behalf Of Zhigang Gong
Sent: Monday, July 22, 2013 11:26 AM
To: Edward Ching
Cc: beignet at lists.freedesktop.org
Subject: Re: [Beignet] compiler asserts on unsupported instrinsics compiling kernel with a single function

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

_______________________________________________
Beignet mailing list
Beignet at lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet