<html xmlns:v="urn:schemas-microsoft-com:vml" xmlns:o="urn:schemas-microsoft-com:office:office" xmlns:w="urn:schemas-microsoft-com:office:word" xmlns:m="http://schemas.microsoft.com/office/2004/12/omml" xmlns="http://www.w3.org/TR/REC-html40"><head><meta http-equiv=Content-Type content="text/html; charset=us-ascii"><meta name=Generator content="Microsoft Word 15 (filtered medium)"><style></style></head><body lang=ZH-CN link="#0563C1" vlink="#954F72"><div class=WordSection1><p class=MsoNormal><span lang=EN-US style='font-size:10.5pt;font-family:"Calibri","sans-serif";color:#1F497D'>Hi Edward,<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:10.5pt;font-family:"Calibri","sans-serif";color:#1F497D'><o:p> </o:p></span></p><div style='border:none;border-left:solid blue 1.5pt;padding:0cm 0cm 0cm 4.0pt'><div><div style='border:none;border-top:solid #E1E1E1 1.0pt;padding:3.0pt 0cm 0cm 0cm'><p class=MsoNormal><b><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri","sans-serif"'>From:</span></b><span lang=EN-US style='font-size:11.0pt;font-family:"Calibri","sans-serif"'> beignet-bounces+zhigang.gong=linux.intel.com@lists.freedesktop.org [mailto:beignet-bounces+zhigang.gong=linux.intel.com@lists.freedesktop.org] <b>On Behalf Of </b>Edward Ching<br><b>Sent:</b> Thursday, June 06, 2013 8:03 AM<br><b>To:</b> beignet@lists.freedesktop.org<br><b>Subject:</b> [Beignet] clEnqueueNDRangeKernel and kernel completion<o:p></o:p></span></p></div></div><p class=MsoNormal><span lang=EN-US><o:p> </o:p></span></p><div><div><div><div><p class=MsoNormal style='margin-bottom:12.0pt'><span lang=EN-US>I hope this is the right forum to post comments/questions on Beignet OpenCL API behaviour. If not, please ignore and excuse the disruption.<o:p></o:p></span></p><p class=MsoNormal style='margin-bottom:12.0pt'><b><i><span lang=EN-US style='font-size:10.5pt;font-family:"Calibri","sans-serif";color:#1F497D'>[Gong, Zhigang] You are right, this is the only official place to discuss anything about Beignet project. You are welcome to post related topics here.</span></i></b><span lang=EN-US style='font-size:10.5pt;font-family:"Calibri","sans-serif";color:#1F497D'><o:p></o:p></span></p></div><p class=MsoNormal style='margin-bottom:12.0pt'><span lang=EN-US>I'm running Beignet on an Ivy Bridge machine and noticed that clFinish would return before the GPU has complete processing of previously submitted commands.<br><br>e.g: <br>I submitted an OpenCL kernel via clEnqueueNDRangeKernel, followed by clFinish, and expected the GPU to have finished all processing when clFinish returns.<br><br>But clFinish returned right away, and when I then call clEnqueueMapBuffer to access data, the call blocks, so I traced the logic all the way into the Ivy Bridge GPU device driver (~/drivers/gpu/drm/i915/*), and it looks like every IvyBr GPU batchbuffer used to submit an OpenCL kernel has an associated sequence number which the GPU would write to a special location which has to be monitored in order to tell if the GPU has finished processing the submtted kernel (details in Intel HD graphics PRM and i915/GEM design notes). Neither clFinish nor clEnqueueNDRangeKernel does this monitoring, and clEnqueueMapBuffer happened to do it when it tried to transfer the buffer objects from GPU to CPU domain.<o:p></o:p></span></p></div><p class=MsoNormal style='margin-bottom:12.0pt'><span lang=EN-US>Does this make sense? It seems to me that clFinish should instead be monitoring and blocking if the GPU is still busy executing an OpenCL kernel.<o:p></o:p></span></p><p class=MsoNormal><b><i><span lang=EN-US style='font-size:10.5pt;font-family:"Calibri","sans-serif";color:#1F497D'>[Gong, Zhigang] </span></i></b><span lang=EN-US style='font-size:10.5pt;font-family:"Calibri","sans-serif";color:#1F497D'>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.<o:p></o:p></span></p><p class=MsoNormal><span lang=EN-US style='font-size:10.5pt;font-family:"Calibri","sans-serif";color:#1F497D'><o:p> </o:p></span></p></div><p class=MsoNormal><span lang=EN-US>/Ed<o:p></o:p></span></p></div></div></div></body></html>