<div dir="ltr">It was from commit 1c3d7709dec22c61d9c3105e4838edce8e6ac014 (HEAD -> master, origin/master, origin/HEAD)<br><div>Date: Sun Jun 14 16:24:18 2020 +0100</div><div><br></div><div>I added those synchronizations for debugging purposes.</div>Ye<br><div><div><div><div dir="ltr" class="gmail_signature" data-smartmail="gmail_signature"><div dir="ltr"><div><div dir="ltr">===================<br>
Ye Luo, Ph.D.<br>Computational Science Division & Leadership Computing Facility<br>
Argonne National Laboratory</div></div></div></div></div><br></div></div></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Tue, Jun 16, 2020 at 8:55 PM Shilei Tian <<a href="mailto:shiltian@cs.stonybrook.edu">shiltian@cs.stonybrook.edu</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"><div style="overflow-wrap: break-word;">Hi Ye,<div><br></div><div>Thanks for your report. From the code snippet, you seem to use a very old version of the RTL. Would you please try the trunk as we don’t have those kind of synchronization.</div><div><br><div>
<div dir="auto" style="text-align:start;text-indent:0px">Regards,<br>Shilei Tian<br>Ph.D. Student<br>Exasca||ab<br>Institute of Advanced Computational Science<br>Stony Brook University</div>
</div>
<div><br><blockquote type="cite"><div>On Jun 16, 2020, at 21:26, Ye Luo via Openmp-dev <<a href="mailto:openmp-dev@lists.llvm.org" target="_blank">openmp-dev@lists.llvm.org</a>> wrote:</div><br><div><div dir="ltr"><div>Hi all,</div><div>I got libomptarget failure when offloading the same kernel function from multiple host threads. I reported the bug as <a href="https://bugs.llvm.org/show_bug.cgi?id=46257" target="_blank">https://bugs.llvm.org/show_bug.cgi?id=46257</a>.</div><div></div><div><br></div><div>After a bit investigation, I'm able to narrow down a bit the scope of the issue.</div><div></div><div>The error from CUDA runtime is "an illegal memory access was encountered".</div><div><br></div><div><div>A) On the libomptarget side<br></div></div><div>In libomptarget/plugins/cuda/src/rtl.cpp, I added a few synchronization before and after cuLaunchKernel and a mutex to protect the kernel execution.<br></div><div><br></div><div> Err = cuCtxSynchronize();<br> if (!checkResult(Err, "Error before cuLaunchKernel\n"))<br> return OFFLOAD_FAIL;<br><br> //kernelmtx.lock();<br> Err = cuLaunchKernel(KernelInfo->Func, CudaBlocksPerGrid, /* gridDimY */ 1,<br> /* gridDimZ */ 1, CudaThreadsPerBlock,<br> /* blockDimY */ 1, /* blockDimZ */ 1,<br> /* sharedMemBytes */ 0, Stream, &Args[0], nullptr);<br> if (!checkResult(Err, "Error returned from cuLaunchKernel\n"))<br> return OFFLOAD_FAIL;<br><br> Err = cuCtxSynchronize();<br> //kernelmtx.unlock();<br> if (!checkResult(Err, "Error after cuLaunchKernel\n"))<br> return OFFLOAD_FAIL;</div><div><br></div><div>Without the mutex. the first error from all the threads is "Error after cuLaunchKernel". After enabling the mutex, the code runs well. If I move the mutex unlock() before the cuCtxSynchronize, the code still runs into error. So I think the error comes from kernel execution not something else.</div><div><br></div><div>2) on my application side, I tried to see what triggers the error. Once I commented out "omp parallel" in the hierarchical parallelism. <a href="https://github.com/QMCPACK/miniqmc/blob/5a11c4131dbf91bf9f9977057a03aee485368f0d/src/QMCWaveFunctions/einspline_spo_omp.cpp#L258" target="_blank">https://github.com/QMCPACK/miniqmc/blob/5a11c4131dbf91bf9f9977057a03aee485368f0d/src/QMCWaveFunctions/einspline_spo_omp.cpp#L258</a><br></div><div>Multi host threads offloading start to run without error.</div><div><br></div><div>I'd like to ask the wider community to see potential causes of this bug.<br></div><div>Is there a thread unsafe state machine with a kernel generated from a target region? Or any other potential reason for the error?</div><div><br></div><div>Best,</div><div>Ye<br></div><div><div dir="ltr"><div dir="ltr"><div><div dir="ltr">===================<br>
Ye Luo, Ph.D.<br>Computational Science Division & Leadership Computing Facility<br>
Argonne National Laboratory</div></div></div></div></div></div>
_______________________________________________<br>Openmp-dev mailing list<br><a href="mailto:Openmp-dev@lists.llvm.org" target="_blank">Openmp-dev@lists.llvm.org</a><br><a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev" target="_blank">https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev</a><br></div></blockquote></div><br></div></div></blockquote></div>