Good idea. I agree to add this option by default currently. Could you make
such a patch?
Thanks.
On Mon, Jul 22, 2013 at 12:08 PM, Yang, Rong R <rong.r.y...@intel.com>wrote:
> Can we simply add clang build option -fno-builtin to disable these
> intrinsic?
>
> -----Original Message-----
> From: beignet-bounces+rong.r.yang=intel....@lists.freedesktop.org [mailto:
> beignet-bounces+rong.r.yang=intel....@lists.freedesktop.org] On Behalf Of
> Zhigang Gong
> Sent: Monday, July 22, 2013 11:26 AM
> To: Edward Ching
> Cc: beignet@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@lists.freedesktop.org
> > http://lists.freedesktop.org/mailman/listinfo/beignet
>
> _______________________________________________
> Beignet mailing list
> Beignet@lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
> _______________________________________________
> Beignet mailing list
> Beignet@lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
>
_______________________________________________
Beignet mailing list
Beignet@lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet
