<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">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">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" 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></div>