<div dir="ltr">Thank you so much for your responses. Here's a simplified test case I used to verify that clFinish would block correctly for exec completion of previously submitted kernel. Not sure if attachements can be sent to this mail list, so I'm cut & pasting the code here (2 files: ~/beignet/utests/compiler_cl_finish.cpp, and ~/beignet/kernels/<a href="http://test_cl_finsh.cl">test_cl_finsh.cl</a>)<div> <br></div><div>/Ed</div><div><br><div><div class="gmail_extra">1) compiler_cl_finish.cpp</div><div class="gmail_extra"><br></div><div class="gmail_extra"><div class="gmail_extra">#include "utest_helper.hpp"</div> <div class="gmail_extra">#include <sys/time.h></div><div class="gmail_extra"><br></div><div class="gmail_extra">#define T_GET(t) gettimeofday(&t, NULL);</div><div class="gmail_extra">#define T_LAPSE(t1, t2) \</div> <div class="gmail_extra"> ((t2.tv_sec+t2.tv_usec*0.000001) - (t1.tv_sec+t1.tv_usec*0.000001))</div><div class="gmail_extra"> </div><div class="gmail_extra">static void compiler_cl_finish(void)</div><div class="gmail_extra"> {</div><div class="gmail_extra"> const size_t n = 16*1024*1024; </div><div class="gmail_extra"> struct timeval t1, t2;</div><div class="gmail_extra"> float t_fin, t_map_w_fin,t_map_wo_fin;</div><div class="gmail_extra"> </div><div class="gmail_extra"> // Setup kernel and buffers</div><div class="gmail_extra"> OCL_CREATE_KERNEL("test_cl_finish");</div><div class="gmail_extra"> OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(int), NULL);</div> <div class="gmail_extra"> OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(int), NULL);</div><div class="gmail_extra"><br></div><div class="gmail_extra"> // Run the kernel</div><div class="gmail_extra"> locals[0] = 64; </div> <div class="gmail_extra"> globals[0] = 32 * locals[0];</div><div class="gmail_extra"> OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);</div><div class="gmail_extra"> OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);</div><div class="gmail_extra"> OCL_SET_ARG(2, sizeof(int), &n);</div><div class="gmail_extra"> OCL_SET_ARG(3, sizeof(int), &globals[0]); </div><div class="gmail_extra"><br></div><div class="gmail_extra"> // 1st time map after clFinish </div> <div class="gmail_extra"> OCL_NDRANGE(1);</div><div class="gmail_extra"> T_GET(t1);</div><div class="gmail_extra"> OCL_FINISH();</div><div class="gmail_extra"> T_GET(t2);</div><div class="gmail_extra"> t_fin = T_LAPSE(t1, t2); </div> <div class="gmail_extra"> </div><div class="gmail_extra"> T_GET(t1); </div><div class="gmail_extra"> OCL_MAP_BUFFER(0);</div><div class="gmail_extra"> T_GET(t2);</div><div class="gmail_extra"> t_map_w_fin = T_LAPSE(t1, t2);</div> <div class="gmail_extra"><br></div><div class="gmail_extra"> // 2nd time map without clFinish</div><div class="gmail_extra"> OCL_NDRANGE(1);</div><div class="gmail_extra"> T_GET(t1); </div><div class="gmail_extra"> OCL_MAP_BUFFER(0);</div> <div class="gmail_extra"> T_GET(t2); </div><div class="gmail_extra"> t_map_wo_fin = T_LAPSE(t1, t2); </div><div class="gmail_extra"> </div><div class="gmail_extra"> OCL_ASSERT(t_fin > t_map_w_fin && t_map_wo_fin > t_map_w_fin); </div> <div class="gmail_extra"> OCL_UNMAP_BUFFER(0);</div><div class="gmail_extra">}</div><div class="gmail_extra"><br></div><div class="gmail_extra">MAKE_UTEST_FROM_FUNCTION(compiler_cl_finish);</div><div class="gmail_extra"> <br></div><div class="gmail_extra" style>2) <a href="http://test_cl_finish.cl">test_cl_finish.cl</a></div><div class="gmail_extra" style><br></div><div class="gmail_extra" style><div class="gmail_extra"><br></div><div class="gmail_extra"> __kernel void</div><div class="gmail_extra">test_cl_finish(__global int *src, __global int *dst, int n, int num_threads)</div><div class="gmail_extra">{</div><div class="gmail_extra"><span class="" style="white-space:pre"> </span>int tid, pos;</div> <div class="gmail_extra"><span class="" style="white-space:pre"> </span></div><div class="gmail_extra"><span class="" style="white-space:pre"> </span>tid = get_global_id(0);</div><div class="gmail_extra"><span class="" style="white-space:pre"> </span>for (pos=tid; pos < n; pos+=num_threads) {</div> <div class="gmail_extra"><span class="" style="white-space:pre"> </span>dst[pos] = src[pos];</div><div class="gmail_extra"><span class="" style="white-space:pre"> </span>}</div><div class="gmail_extra">}</div><div><br></div> </div><br><div class="gmail_quote">On Wed, Jun 5, 2013 at 10:45 PM, Zou, Nanhai <span dir="ltr"><<a href="mailto:nanhai.zou@intel.com" target="_blank">nanhai.zou@intel.com</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left-width:1px;border-left-color:rgb(204,204,204);border-left-style:solid;padding-left:1ex"> <div lang="ZH-CN" link="#0563C1" vlink="#954F72"> <div> <p class=""><span lang="EN-US" style="font-size:10.5pt;font-family:Calibri,sans-serif;color:rgb(31,73,125)"><u></u> <u></u></span></p> <div style="border-style:none none none solid;border-left-color:blue;border-left-width:1.5pt;padding:0cm 0cm 0cm 4pt"> <div><div class="im"> <div> <p class="" style="margin-bottom:12pt"><span lang="EN-US"><u></u> <u></u></span></p> <p class=""><b><i><span lang="EN-US" style="font-size:10.5pt;font-family:Calibri,sans-serif;color:rgb(31,73,125)">>></span></i></b><b><i><span lang="EN-US" style="font-size:10.5pt;font-family:Calibri,sans-serif;color:rgb(31,73,125)">[Gong, Zhigang] </span></i></b><span lang="EN-US" style="font-size:10.5pt;font-family:Calibri,sans-serif;color:rgb(31,73,125)">Your finding is correct, current clFinish does nothing which is not comply with the spec. It is on our TODO list. Actually, we have more related TODO items. Currently, the clEnqueueNDRangeKernel flushes the batchbuffer every time, and thus the clFlush is also an empty function. We also need to optimize it to track the states and avoid some unnecessary pipe controls for each kernel. But our team’s current focus is to implement the missing opencl features, and try hard to pass the piglit test. After that, we will turn to these items. And as usual, if everyone from the community want to contribute on these items, we will be more than happy to review and accept it.<u></u><u></u></span></p> </div> <p class=""><span lang="EN-US" style="color:rgb(31,73,125)"><u></u> <u></u></span></p> </div><p class=""><span lang="EN-US" style="font-size:10.5pt;font-family:Calibri,sans-serif;color:rgb(31,73,125)">Unlike the cl_event and asynchronize things in our plan, this clFinish one looks like a bug. <u></u><u></u></span></p> <p class=""><span lang="EN-US" style="font-size:10.5pt;font-family:Calibri,sans-serif;color:rgb(31,73,125)">Code is there, just not work as expected.<u></u><u></u></span></p> <p class=""><span lang="EN-US" style="font-size:10.5pt;font-family:Calibri,sans-serif;color:rgb(31,73,125)"><u></u> <u></u></span></p> <p class=""><span lang="EN-US" style="font-size:10.5pt;font-family:Calibri,sans-serif;color:rgb(31,73,125)">Edward, would you mind provide a simple test case for this?<u></u><u></u></span></p> <p class=""><span lang="EN-US" style="font-size:10.5pt;font-family:Calibri,sans-serif;color:rgb(31,73,125)"><u></u> <u></u></span></p> <p class=""><span lang="EN-US" style="font-size:10.5pt;font-family:Calibri,sans-serif;color:rgb(31,73,125)">Thanks<span class=""><font color="#888888"><u></u><u></u></font></span></span></p><span class=""><font color="#888888"> <p class=""><span lang="EN-US" style="font-size:10.5pt;font-family:Calibri,sans-serif;color:rgb(31,73,125)">Zou Nanhai<u></u><u></u></span></p> <p class=""><span lang="EN-US" style="font-size:10.5pt;font-family:Calibri,sans-serif;color:rgb(31,73,125)"><u></u> <u></u></span></p> </font></span></div> </div> </div> </div> </blockquote></div><br></div></div></div></div>