[Openmp-dev] Potential offload kernel data race
Ye Luo via Openmp-dev
openmp-dev at lists.llvm.org
Thu Jun 25 12:05:47 PDT 2020
I cannot use openmp-cuda-mode. I don't remember now exactly where
(reduction?) but it breaks other things in my code.
Ye
===================
Ye Luo, Ph.D.
Computational Science Division & Leadership Computing Facility
Argonne National Laboratory
On Thu, Jun 25, 2020 at 2:02 PM Alexey.Bataev <a.bataev at outlook.com> wrote:
> You can use shared memory directly by using #pragma omp allocate and
> allocate clauses. Plus, shared memory is actively used by the runtime. But
> I agree, that it would be good to allow to control the size.
>
> BTW, you can disable this globalization mechanism by using
> -fopenmp-cuda-mode option. In this case, you need to control sharing the
> variables yourself.
>
> -------------
> Best regards,
> Alexey Bataev
>
> 25.06.2020 2:51 PM, Ye Luo пишет:
>
> I think hard-coded size 128 is suboptimal. Can we allow users to select
> that?
> Since we cannot use shared memory directly, I believe there are plenty of
> shared memory available.
> It seems that my test code is exceeding 128 byte and I'd like to set a
> higher limit.
> Ye
> ===================
> Ye Luo, Ph.D.
> Computational Science Division & Leadership Computing Facility
> Argonne National Laboratory
>
>
> On Thu, Jun 25, 2020 at 1:02 PM Alexey Bataev <a.bataev at hotmail.com>
> wrote:
>
>> Globalized variables are the variables that must be shared between the
>> threads but initially allocated by the master thread. On GPU, local memory
>> is private for threads, and the local variables cannkt be shared between
>> the threads. But OpenMP requires that such variables must be shared between
>> the threads. So, if we see, that the local variable is allocated in the
>> context of the master thread and later is shared in parallel region (or
>> potentially can be shared, if it may escape its declaration context by the
>> reference or by the address), such variables are allocated either in shared
>> or in global memory. If the total size of globalized variables is less than
>> 128 bytes, they are allocated in shared memory. Otherwise, fallback to
>> malloc. If -fno-penmp-cuda-parallel-target-region is used, then again, if
>> total size <= 128 bytes - shared memory is used, otherwise statically
>> allocated global memory buffer is used.
>>
>> Best regards,
>> Alexey Bataev
>>
>> 25 июня 2020 г., в 13:45, Ye Luo <xw111luoye at gmail.com> написал(а):
>>
>>
>> Could you add documentation for this option?
>> fopenmp-cuda-parallel-target-region
>> The name doesn't reflect the actual effect.
>> The differential review mentions
>> ```
>> Added support for dynamic memory allocation for globalized variables in
>> case if execution of target regions in parallel is required.
>> ```
>> So what are the globalized variables? When do they occur? since it
>> impacts performance, application developers desire to know more details.
>> ```
>>
>> It might use "slow" allocation functions, in general, since it may use
>> malloc on the device side. Being disabled, it uses statically preallocated
>> memory, which might be faster, if parallel target regions are not required.
>>
>> ```
>> In the slow code path, does it always call malloc in the device or malloc
>> gets called when exceeding a size limit? Is the call explicitly generated
>> by clang or the cuda driver handles it?
>> I saw some UseSharedMemory flags in the review. Is the slow code path
>> uses CUDA shared memory and it may need global memory when the size limit
>> is exceeded?
>> In the fast code path, what is the scope of this statically preallocated
>> memory? Is it owned by the CUDA plugin?
>>
>> Thanks,
>> Ye
>> ===================
>> Ye Luo, Ph.D.
>> Computational Science Division & Leadership Computing Facility
>> Argonne National Laboratory
>>
>>
>> On Thu, Jun 25, 2020 at 9:44 AM Alexey.Bataev <a.bataev at outlook.com>
>> wrote:
>>
>>> Hi Ye, the patch is committed, you can check if the compiler works for
>>> you.
>>>
>>> -------------
>>> Best regards,
>>> Alexey Bataev
>>>
>>> 22.06.2020 1:29 PM, Ye Luo via Openmp-dev пишет:
>>>
>>> It is working well now. I tested both the miniapp and the full
>>> application.
>>> ===================
>>> Ye Luo, Ph.D.
>>> Computational Science Division & Leadership Computing Facility
>>> Argonne National Laboratory
>>>
>>>
>>> On Mon, Jun 22, 2020 at 11:56 AM Alexey.Bataev <a.bataev at outlook.com>
>>> wrote:
>>>
>>>> Hi Ye, could you try your code with the attached patch for the clang?
>>>>
>>>> -------------
>>>> Best regards,
>>>> Alexey Bataev
>>>>
>>>> 19.06.2020 6:45 PM, Ye Luo via Openmp-dev пишет:
>>>>
>>>> Yes. It makes the code working.
>>>> Do you have any idea about the code pattern?
>>>> Before the compiler fix gets in, I can put the workaround in the real
>>>> code which has multiple places similar to the miniapp.
>>>> Best,
>>>> Ye
>>>> ===================
>>>> Ye Luo, Ph.D.
>>>> Computational Science Division & Leadership Computing Facility
>>>> Argonne National Laboratory
>>>>
>>>>
>>>> On Fri, Jun 19, 2020 at 4:35 PM Alexey.Bataev <a.bataev at outlook.com>
>>>> wrote:
>>>>
>>>>> Try to use `omp parallel firstprivate(first, last,
>>>>> offload_scratch_iVP_ptr, psiinv_ptr, i)` as the inner parallel directive.
>>>>>
>>>>> -------------
>>>>> Best regards,
>>>>> Alexey Bataev
>>>>>
>>>>> 16.06.2020 9:26 PM, Ye Luo via Openmp-dev пишет:
>>>>>
>>>>> Hi all,
>>>>> I got libomptarget failure when offloading the same kernel function
>>>>> from multiple host threads. I reported the bug as
>>>>> https://bugs.llvm.org/show_bug.cgi?id=46257.
>>>>>
>>>>> After a bit investigation, I'm able to narrow down a bit the scope of
>>>>> the issue.
>>>>> The error from CUDA runtime is "an illegal memory access was
>>>>> encountered".
>>>>>
>>>>> A) On the libomptarget side
>>>>> In libomptarget/plugins/cuda/src/rtl.cpp, I added a few
>>>>> synchronization before and after cuLaunchKernel and a mutex to protect the
>>>>> kernel execution.
>>>>>
>>>>> Err = cuCtxSynchronize();
>>>>> if (!checkResult(Err, "Error before cuLaunchKernel\n"))
>>>>> return OFFLOAD_FAIL;
>>>>>
>>>>> //kernelmtx.lock();
>>>>> Err = cuLaunchKernel(KernelInfo->Func, CudaBlocksPerGrid, /*
>>>>> gridDimY */ 1,
>>>>> /* gridDimZ */ 1, CudaThreadsPerBlock,
>>>>> /* blockDimY */ 1, /* blockDimZ */ 1,
>>>>> /* sharedMemBytes */ 0, Stream, &Args[0],
>>>>> nullptr);
>>>>> if (!checkResult(Err, "Error returned from cuLaunchKernel\n"))
>>>>> return OFFLOAD_FAIL;
>>>>>
>>>>> Err = cuCtxSynchronize();
>>>>> //kernelmtx.unlock();
>>>>> if (!checkResult(Err, "Error after cuLaunchKernel\n"))
>>>>> return OFFLOAD_FAIL;
>>>>>
>>>>> 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.
>>>>>
>>>>> 2) on my application side, I tried to see what triggers the error.
>>>>> Once I commented out "omp parallel" in the hierarchical parallelism.
>>>>> https://github.com/QMCPACK/miniqmc/blob/5a11c4131dbf91bf9f9977057a03aee485368f0d/src/QMCWaveFunctions/einspline_spo_omp.cpp#L258
>>>>> Multi host threads offloading start to run without error.
>>>>>
>>>>> I'd like to ask the wider community to see potential causes of this
>>>>> bug.
>>>>> Is there a thread unsafe state machine with a kernel generated from a
>>>>> target region? Or any other potential reason for the error?
>>>>>
>>>>> Best,
>>>>> Ye
>>>>> ===================
>>>>> Ye Luo, Ph.D.
>>>>> Computational Science Division & Leadership Computing Facility
>>>>> Argonne National Laboratory
>>>>>
>>>>> _______________________________________________
>>>>> Openmp-dev mailing listOpenmp-dev at lists.llvm.orghttps://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev
>>>>>
>>>>>
>>>> _______________________________________________
>>>> Openmp-dev mailing listOpenmp-dev at lists.llvm.orghttps://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev
>>>>
>>>>
>>> _______________________________________________
>>> Openmp-dev mailing listOpenmp-dev at lists.llvm.orghttps://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev
>>>
>>> _______________________________________________
>> Openmp-dev mailing list
>> Openmp-dev at lists.llvm.org
>> https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev
>>
>>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/openmp-dev/attachments/20200625/96b282a0/attachment-0001.html>
More information about the Openmp-dev
mailing list