<div dir="ltr"><div>Good idea. I agree to add this option by default currently. Could you make such a patch?</div><div>Thanks.</div></div><div class="gmail_extra"> <div class="gmail_quote">On Mon, Jul 22, 2013 at 12:08 PM, Yang, Rong R <<a href="mailto:rong.r.yang@intel.com" target="_blank">rong.r.yang@intel.com</a>> wrote: <blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Can we simply add clang build option -fno-builtin to disable these intrinsic? <div class="HOEnZb"><div class="h5"> -----Original Message----- From: beignet-bounces+rong.r.yang=<a href="mailto:intel.com@lists.freedesktop.org">intel.com@lists.freedesktop.org</a> [mailto:<a href="mailto:beignet-bounces%2Brong.r.yang">beignet-bounces+rong.r.yang</a>=<a href="mailto:intel.com@lists.freedesktop.org">intel.com@lists.freedesktop.org</a>] On Behalf Of Zhigang Gong Sent: Monday, July 22, 2013 11:26 AM To: Edward Ching Cc: <a href="mailto:beignet@lists.freedesktop.org">beignet@lists.freedesktop.org</a> 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 > > <a href="http://test_split.cl" target="_blank">test_split.cl</a>: > > __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 > <a href="mailto:Beignet@lists.freedesktop.org">Beignet@lists.freedesktop.org</a> > <a href="http://lists.freedesktop.org/mailman/listinfo/beignet" target="_blank">http://lists.freedesktop.org/mailman/listinfo/beignet</a> _______________________________________________ Beignet mailing list <a href="mailto:Beignet@lists.freedesktop.org">Beignet@lists.freedesktop.org</a> <a href="http://lists.freedesktop.org/mailman/listinfo/beignet" target="_blank">http://lists.freedesktop.org/mailman/listinfo/beignet</a> _______________________________________________ Beignet mailing list <a href="mailto:Beignet@lists.freedesktop.org">Beignet@lists.freedesktop.org</a> <a href="http://lists.freedesktop.org/mailman/listinfo/beignet" target="_blank">http://lists.freedesktop.org/mailman/listinfo/beignet</a> </div></div></blockquote></div> </div>