<div dir="ltr"><div><div>It compiles fine now with the latest master branch with the tip at <br>commit 54d31547b9e4ce10e58a6e116ee33b5aba7b4ab8<br>Author: Yi Sun <<a href="mailto:yi.sun@intel.com">yi.sun@intel.com</a>><br>
Date:   Mon Jul 22 15:59:42 2013 +0800<br>utest: add built-in test case for get_global_id.<br>    <br></div>Thanks!<br></div>/Ed<br><div><div><br></div></div></div><div class="gmail_extra"><br><br><div class="gmail_quote">
On Mon, Jul 22, 2013 at 12:29 AM, Zhigang Gong <span dir="ltr"><<a href="mailto:zhigang.gong@linux.intel.com" target="_blank">zhigang.gong@linux.intel.com</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
Hi Edward,<br>
<br>
Could you try to compile your kernel again with latest master branch?<br>
Rong's patch should get this issue fixed.<br>
<div class="im HOEnZb"><br>
On Sat, Jul 20, 2013 at 10:49:08PM -0700, Edward Ching wrote:<br>
</div><div class="HOEnZb"><div class="h5">> When compiling a simple single function kernel (patch attached, also cut<br>
> and pasted below), the compiler asserted on unsupported intrinsics:<br>
><br>
> ASSERTION FAILED: Unsupported intrinsics<br>
>   at file<br>
> /root/WORK/test_split/beignet/backend/src/llvm/llvm_gen_backend.cpp,<br>
> function void gbe::GenWriter::regAllocateCallInst(llvm::CallInst&), line<br>
> 1688<br>
> Stack dump:<br>
> 0. Running pass 'Function Pass Manager' on module '/tmp/file0yciwa.ll'.<br>
> 1. Running pass 'Gen Back-End' on function '@test_split'<br>
><br>
> It looks like this is triggered by the "shr_mem[pid+offset]=0;" line when 0<br>
> is assigned. Is this a bug? It looks like the kernel code is not doing<br>
> anything illegal. It compiled ok using Intel OpenCL SDK 2013 (on Linux, for<br>
> CPU)<br>
><br>
> Thanks,<br>
> /Ed<br>
><br>
> <a href="http://test_split.cl" target="_blank">test_split.cl</a>:<br>
><br>
> __kernel void<br>
> test_split(<br>
> __global int *in,<br>
> int stride,<br>
> int n_recs,<br>
> int n_parts,<br>
> __global int *out,<br>
> __local  int *shr_mem)<br>
> {<br>
> int pid;<br>
> int glb_tid = get_global_id(0);<br>
> int loc_tid = get_local_id(0);<br>
> int offset  = loc_tid*n_parts;<br>
><br>
> for(pid=0; pid<n_parts; pid++)<br>
> {<br>
> shr_mem[pid+offset]=0;<br>
> }<br>
><br>
> for(int pos=glb_tid; pos<n_recs; pos+=stride)<br>
> {<br>
> pid=in[pos];<br>
> shr_mem[pid+offset]++;<br>
> }<br>
><br>
> for(pid=0; pid<n_parts; pid++)<br>
> {<br>
> out[glb_tid + stride*pid] = shr_mem[pid+offset];<br>
> }<br>
> }<br>
><br>
> compiler_test_split.cpp:<br>
><br>
> void compiler_test_split(void)<br>
> {<br>
>   OCL_CREATE_KERNEL("test_split");<br>
> }<br>
<br>
<br>
</div></div><div class="HOEnZb"><div class="h5">> _______________________________________________<br>
> Beignet mailing list<br>
> <a href="mailto:Beignet@lists.freedesktop.org">Beignet@lists.freedesktop.org</a><br>
> <a href="http://lists.freedesktop.org/mailman/listinfo/beignet" target="_blank">http://lists.freedesktop.org/mailman/listinfo/beignet</a><br>
<br>
</div></div></blockquote></div><br></div>