<div dir="ltr">Hi Xiuli,<br><div><br></div><div>Apologies if I were not clear enough with my questions.</div><div><br></div><div>1. By saying that there is only 64K local memory, I assume you meant to hint that this is a scarce resource. We use only ~2KB local memory per work group required mostly for prefetching and reduction across work-group. By disabling pre-fetching we could get this down to a minimum of about 800 bytes. However, do you expect that to help in any way? Is the hardware capable of keeping in flight >30-32 waves of 64 threads?</div><div><br></div>Also, I think I'm lacking some detailed knowledge as I do not see how is this related to the drm_intel_gem_bo_context_exec() issue.<div><br></div><div>2. As mentioned above, I have local work size = 64 and rely on splitting the work over the global grid (so small workloads will have 100s, large ones 10000s larger global work size).</div><div><br></div><div>3. Do you mean that I should *not* expect concurrency between CPU and GPU to be possible with beignet and clEnqueueWriteBuffer, clEnqueuereadBuffer, and clEnqueueNDRangeKernel will exhibit blocking behavior?</div><div><br></div><div><br></div><div>I have not had time to file a bugzilla yet with reproduction details, sorry about that. </div><div>The GROMACS source you'll need is in the master branch, plus the <a href="https://gerrit.gromacs.org/#/c/5752/2">https://gerrit.gromacs.org/#/c/5752/2</a> change under review fixes some execution width assumptions. Other than that you'll need a small patch to enable Intel iGPUs (e.g found here <a href="https://bugs.freedesktop.org/show_bug.cgi?id=94265">https://bugs.freedesktop.org/show_bug.cgi?id=94265</a> which is BTW another bug on IVB). You'll also run into the include path issue I mentioned before for which you'll need to activate the Apple workaround, here's a patch:</div><div><br></div><div><div>diff --git a/src/gromacs/gpu_utils/gpu_utils_ocl.cpp b/src/gromacs/gpu_utils/gpu_utils_ocl.cpp</div><div>index 2084d8c..8928582 100644</div><div>--- a/src/gromacs/gpu_utils/gpu_utils_ocl.cpp</div><div>+++ b/src/gromacs/gpu_utils/gpu_utils_ocl.cpp</div><div>@@ -131,6 +131,8 @@ static int is_gmx_supported_gpu_id(struct gmx_device_info_t *ocl_gpu_device)</div><div> return egpuCompatible;</div><div> case OCL_VENDOR_AMD:</div><div> return runningOnCompatibleOSForAmd() ? egpuCompatible : egpuIncompatible;</div><div>+ case OCL_VENDOR_INTEL:</div><div>+ return egpuCompatible;</div><div> default:</div><div> return egpuIncompatible;</div><div> }</div><div>diff --git a/src/gromacs/gpu_utils/ocl_compiler.cpp b/src/gromacs/gpu_utils/ocl_compiler.cpp</div><div>index 6a4772a..9aa3c1e 100644</div><div>--- a/src/gromacs/gpu_utils/ocl_compiler.cpp</div><div>+++ b/src/gromacs/gpu_utils/ocl_compiler.cpp</div><div>@@ -747,7 +747,8 @@ ocl_get_build_options_string(cl_context context,</div><div> * OpenCL implementations are happy with. Since the standard still says</div><div> * it should be quoted, we handle Apple as a special case.</div><div> */</div><div>-#ifdef __APPLE__</div><div>+//#ifdef __APPLE__</div><div>+#if 1</div><div> std::string unescaped_ocl_root_path = get_ocl_root_path();</div><div> std::string ocl_root_path;</div></div><div><br></div><div><br><div>Additionally, here's an input file you'll need to be able to start the program:</div><div><a href="https://www.dropbox.com/s/hm5t90iwo3xw5ws/water-48k-frozen.tpr?dl=0">https://www.dropbox.com/s/hm5t90iwo3xw5ws/water-48k-frozen.tpr?dl=0</a><br></div><div>which you can do with the following command:<br></div><div>/PATH/gmx mdrun -s water-48k-frozen</div><div><br></div></div><div>Let me know if something is unclear.</div><div><br></div><div>Thanks for the help!</div><div><br></div><div>Cheers,</div></div><div class="gmail_extra"><br clear="all"><div><div class="gmail_signature">--<br>Szilárd</div></div> <br><div class="gmail_quote">On Thu, Mar 31, 2016 at 5:34 AM, Pan, Xiuli <span dir="ltr"><<a href="mailto:xiuli.pan@intel.com" target="_blank">xiuli.pan@intel.com</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"> <div lang="EN-US" link="blue" vlink="purple"> <div> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d">Hi Szilárd,<u></u><u></u></span></p> <p class="MsoNormal"><a name="m_5434114372012745357__MailEndCompose"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d"><u></u> <u></u></span></a></p> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d">Since you have some questions and I could not reproduce them here I just make some response that I think may related to this problem:<u></u><u></u></span></p> <p><u></u><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d"><span>1.<span style="font:7.0pt "Times New Roman""> </span></span></span><u></u><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d">We only have 64K Share local memory for all of the work groups<u></u><u></u></span></p> <p><u></u><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d"><span>2.<span style="font:7.0pt "Times New Roman""> </span></span></span><u></u><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d">The drm_intel_gem_bo_context_exec() failed have a lot of reasons, could give us the detail about your test about the execution wide?<u></u><u></u></span></p> <p><u></u><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d"><span>3.<span style="font:7.0pt "Times New Roman""> </span></span></span><u></u><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d">As far as I know most enqueue in beignet default to be blocking(some related to GPU is not blocking) , you can see that api clFlush is actually an empty function.<u></u><u></u></span></p> <p class="MsoNormal" style="margin-left:.25in"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d"><u></u> <u></u></span></p> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d">Also I am trying to reproduce your bug here and I am setting up GROMACS. Is there anything I should know to run it with beignet?<u></u><u></u></span></p> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d"><u></u> <u></u></span></p> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d">Thanks<u></u><u></u></span></p> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d">Xiuli <u></u><u></u></span></p> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d"><u></u> <u></u></span></p> <p class="MsoNormal"><a name="m_5434114372012745357______replyseparator"></a><b><span style="font-size:11.0pt;font-family:"Calibri",sans-serif">From:</span></b><span style="font-size:11.0pt;font-family:"Calibri",sans-serif"> Beignet [mailto:<a href="mailto:beignet-bounces@lists.freedesktop.org" target="_blank">beignet-bounces@lists.freedesktop.org</a>] <b>On Behalf Of </b>Szilárd Páll<br> <b>Sent:</b> Thursday, March 31, 2016 3:14 AM<br> <b>To:</b> <a href="mailto:beignet@lists.freedesktop.org" target="_blank">beignet@lists.freedesktop.org</a><br> <b>Subject:</b> Re: [Beignet] GROMACS on beignet<u></u><u></u></span></p><div><div class="h5"> <p class="MsoNormal"><u></u> <u></u></p> <div> <p class="MsoNormal">Hello again,<u></u><u></u></p> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal">I have been trying to verify whether there may be assumptions >=32-wide execution hiding in the kernels (in particular in code that's using local memory for prefetching or reduction) and tried dropping in mem fences to test a few things, but at several points I managed to trigger the aforementioned error:<u></u><u></u></p> </div> <div> <div> <p class="MsoNormal">drm_intel_gem_bo_context_exec() failed: Input/output error<u></u><u></u></p> </div> </div> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal">Is this a known issues? There have been reports of it, but perhaps it is just the manifestation of multiple possible issues?<u></u><u></u></p> </div> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal">Secondly, I do not see the reason why I get blocking behavior of all enqueue operations (and I don't get this on NVIDIA or AMD). Are there any peculiarities I should be aware of?<u></u><u></u></p> </div> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal">Cheers,<u></u><u></u></p> </div> </div> <div> <p class="MsoNormal"><br clear="all"> <u></u><u></u></p> <div> <div> <p class="MsoNormal">--<br> Szilárd<u></u><u></u></p> </div> </div> <p class="MsoNormal"><u></u> <u></u></p> <div> <p class="MsoNormal">On Mon, Mar 28, 2016 at 1:49 AM, Szilárd Páll <<a href="mailto:sin.pecado@gmail.com" target="_blank">sin.pecado@gmail.com</a>> wrote:<u></u><u></u></p> <blockquote style="border:none;border-left:solid #cccccc 1.0pt;padding:0in 0in 0in 6.0pt;margin-left:4.8pt;margin-top:5.0pt;margin-right:0in;margin-bottom:5.0pt"> <div> <p class="MsoNormal">Hi Xiuli,<u></u><u></u></p> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal">Thanks for the quick reply!<u></u><u></u></p> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <div> <p class="MsoNormal">On Fri, Mar 25, 2016 at 4:06 AM, Pan, Xiuli <<a href="mailto:xiuli.pan@intel.com" target="_blank">xiuli.pan@intel.com</a>> wrote:<u></u><u></u></p> <blockquote style="border:none;border-left:solid #cccccc 1.0pt;padding:0in 0in 0in 6.0pt;margin-left:4.8pt;margin-top:5.0pt;margin-right:0in;margin-bottom:5.0pt"> <div> <div> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d">Hi Szilárd,</span><u></u><u></u></p> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d"> </span><u></u><u></u></p> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d">What do you mean about quoted includes?</span><u></u><u></u></p> </div> </div> </blockquote> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal">I mean -I"/path/to/headers" does not work, but -I/path/to/headers does.<u></u><u></u></p> </div> <div> <p class="MsoNormal"> <u></u><u></u></p> </div> <blockquote style="border:none;border-left:solid #cccccc 1.0pt;padding:0in 0in 0in 6.0pt;margin-left:4.8pt;margin-top:5.0pt;margin-right:0in;margin-bottom:5.0pt"> <div> <div> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d">If you mean the include in kernels, I think we may have some problem with that. The *.cl we used for clang actually was a copied tmp version stored not in where is used to be. So I think if you just put what need to be included in the old place, clang could not find it. You could try a workaround to pass “-I where/your/header/is” as a build option to clBuildProgram.</span><u></u><u></u></p> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d"> </span><u></u><u></u></p> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d">Then if you have some double types used on Haswell it may have some problem. The hardware for HSW does not support double very well as we have refined our double support to hardware then, so HSW may have some issues with double type. If it is not the problem with double float, you can send your kernel as an attachment or report a bug on our Bugzilla(</span><a href="https://bugs.freedesktop.org" target="_blank"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif">https://bugs.freedesktop.org</span></a><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d">) and we will tried to fix it.</span><u></u><u></u></p> </div> </div> </blockquote> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal">No double precision in the kernels. <u></u><u></u></p> </div> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal">For now I'll post here, I feel like a bug report may be an overkill - especially as I can't provide a full repro case that does not involve building the entire application.<u></u><u></u></p> </div> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal">I've attached a minimum set of source files that's needed to compile. We have pretty heavy preprocessor use that generates kernels for the different inputs / outputs / computation combinations, so one particular flavor that's known to produce incorrect results is generated compiling <a href="http://nbnxn_ocl_kernels.cl" target="_blank"> nbnxn_ocl_kernels.cl</a> with the following flags:<u></u><u></u></p> </div> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal">-D_WARPLESS_SOURCE_ -DGMX_OCL_FASTGEN -DEL_RF -DEELNAME=_ElecRF -DLJ_COMB_GEOM -DVDWNAME=_VdwLJCombGeom -DCENTRAL=22 -DNBNXN_GPU_NCLUSTER_PER_SUPERCLUSTER=8 -DNBNXN_GPU_CLUSTER_SIZE=8 -DNBNXN_GPU_JGROUP_SIZE=4 -DNBNXN_AVOID_SING_R2_INC=1.0e-12f<u></u><u></u></p> </div> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal">Additionally I had a closer look and so far I have observed three issues (additional to the minor include issue mentione before):<u></u><u></u></p> </div> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal">1. If I do a manual prefetch into local memory followed by a mem fence (seenbnxn_ocl_kernel_nowarp.clh line 339), I get the following error:<u></u><u></u></p> </div> <div> <p class="MsoNormal">drm_intel_gem_bo_context_exec() failed: Input/output error<u></u><u></u></p> </div> <div> <p class="MsoNormal">The next kernel call then fails with CL_OUT_OF_RESOURCES.<u></u><u></u></p> </div> <div> <p class="MsoNormal">Without the manual prefetch it works better, but...<u></u><u></u></p> </div> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal">2. The results produced by the kernel are still somewhat off. It could be that I missed a subtle detail and the kernels still do not conform to the hardware's execution model. I'm very familar with Intel's hardware and these kernels were originally designed for 32/64 wide execution.<u></u><u></u></p> </div> </div> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal">3. All task enqueue calls seem to be blocking. <u></u><u></u></p> </div> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <div> <p class="MsoNormal">Thanks & Cheers,<u></u><u></u></p> </div> <div> <p class="MsoNormal">--<u></u><u></u></p> </div> <div> <p class="MsoNormal">Szilárd<u></u><u></u></p> </div> <div> <p class="MsoNormal"><u></u> <u></u></p> </div> <blockquote style="border:none;border-left:solid #cccccc 1.0pt;padding:0in 0in 0in 6.0pt;margin-left:4.8pt;margin-top:5.0pt;margin-right:0in;margin-bottom:5.0pt"> <div> <div> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d"> </span><u></u><u></u></p> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d">Thanks</span><u></u><u></u></p> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d">Xiuli</span><u></u><u></u></p> <p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif;color:#1f497d"> </span><u></u><u></u></p> <p class="MsoNormal"><a name="m_5434114372012745357_m_8161337532417221177_m_-165604280253327"></a><b><span style="font-size:11.0pt;font-family:"Calibri",sans-serif">From:</span></b><span style="font-size:11.0pt;font-family:"Calibri",sans-serif"> Beignet [mailto:</span><a href="mailto:beignet-bounces@lists.freedesktop.org" target="_blank"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif">beignet-bounces@lists.freedesktop.org</span></a><span style="font-size:11.0pt;font-family:"Calibri",sans-serif">] <b>On Behalf Of </b>Szilárd Páll<br> <b>Sent:</b> Friday, March 25, 2016 7:16 AM<br> <b>To:</b> </span><a href="mailto:beignet@lists.freedesktop.org" target="_blank"><span style="font-size:11.0pt;font-family:"Calibri",sans-serif">beignet@lists.freedesktop.org</span></a><span style="font-size:11.0pt;font-family:"Calibri",sans-serif"><br> <b>Subject:</b> [Beignet] GROMACS on beignet</span><u></u><u></u></p> <div> <div> <p class="MsoNormal"> <u></u><u></u></p> <div> <p class="MsoNormal">Hi,<u></u><u></u></p> <div> <p class="MsoNormal"> <u></u><u></u></p> </div> <div> <div> <p class="MsoNormal">I am a developer of the GROMACS (<a href="http://www.gromacs.org" target="_blank">www.gromacs.org</a>) molecular dynamics simulation package. We have OpenCL offload for some of the compute-intensive kernels which that works very well on AMD. I wanted to assess how feasible is to use an Intel iGPU in GROMACS and after jumping through some hoops I got a 4.2 kernel and beignet master installed.<u></u><u></u></p> </div> <div> <p class="MsoNormal"> <u></u><u></u></p> </div> <div> <p class="MsoNormal">Then I ran into the first minor issue: it seems that beignet does not accept quoted includes although AFAIK the double-quoted include paths should be accepted, but that did not work. No big deal, it doesn't work with Apple's OpenCL either, but I thought I'd ask.<u></u><u></u></p> </div> <div> <p class="MsoNormal"> <u></u><u></u></p> </div> <div> <p class="MsoNormal">However, the bigger issue is that running on Haswell (HD 4600, I think) the kernel produces results that are very off (while the very same source gives correct results on other platforms). I've not much time to dig deeper, but I thought I'd drop a mail maybe somebody is interested in helping out with tips or even tracking down where the issue is.<u></u><u></u></p> </div> <div> <p class="MsoNormal"> <u></u><u></u></p> </div> <div> <p class="MsoNormal">Suggestions would be welcome!<u></u><u></u></p> </div> <div> <p class="MsoNormal"> <u></u><u></u></p> </div> <div> <p class="MsoNormal">Cheers,<br clear="all"> <u></u><u></u></p> <div> <div> <p class="MsoNormal">--<br> Szilárd<u></u><u></u></p> </div> </div> </div> </div> </div> </div> </div> </div> </div> </blockquote> </div> <p class="MsoNormal"><u></u> <u></u></p> </div> </div> </div> </blockquote> </div> <p class="MsoNormal"><u></u> <u></u></p> </div> </div></div></div> </div> </blockquote></div><br></div>