[Openmp-dev] Libomptarget fatal error 1: '#pragma omp requires unified_shared_memory' not used consistently!
Joachim Protze via Openmp-dev
openmp-dev at lists.llvm.org
Mon Mar 1 23:47:16 PST 2021
No!
What you suggest (Itaru to check that the system supports "whatever?!")
would be necessary for a fictitious:
#pragma omp assume unified_shared_memory
but certainly not for
#pragma omp require unified_shared_memory
The OpenMP implementation must support the required property or must
state that it is not supported. Causing an error during the execution of
a target region is no compliant behavior.
As Johannes suggested, managed allocators might allow to provide
software USM, when hardware USM is not available. In any case, compiler
and runtime together need to make sure that the requirements for
supporting USM are met or abort otherwise.
BTW: will the managed allocators also change the behavior for stack
allocations? This would be necessary for a compliant USM support.
- Joachim
Am 01.03.21 um 23:49 schrieb Ye Luo via Openmp-dev:
> Are you sure your machine/OS supports USM? Summit like nodes P9+V100 are
> the only architecture with NVIDIA GPU which supports USM that I'm aware of.
As you can read on the page I linked earlier, already Kepler supports
USM. Just not in hardware.
> Ye
> ===================
> Ye Luo, Ph.D.
> Computational Science Division & Leadership Computing Facility
> Argonne National Laboratory
>
>
> On Mon, Mar 1, 2021 at 4:35 PM Itaru Kitayama via Openmp-dev <
> openmp-dev at lists.llvm.org> wrote:
>
>> I’m on JURECA and some nodes are attached to A100 GPUs.
>>
>> On Tue, Mar 2, 2021 at 7:34 Itaru Kitayama <itaru.kitayama at gmail.com>
>> wrote:
>>
>>> Hi all,
>>> In the mean time, what do I do?
>>>
>>> On Tue, Mar 2, 2021 at 3:23 Johannes Doerfert <johannesdoerfert at gmail.com>
>>> wrote:
>>>
>>>> I think that is it. I heard of problems with our USM before.
>>>> We need to use the managed allocators if USM is active, they are
>>>> about to be upstreamed (I hope).
>>>>
>>>>
>>>> On 3/1/21 12:15 PM, Alexey.Bataev wrote:
>>>>> Looks like this example is for Explicit USM and I assume if you
>>>> allocate
>>>>> the memory for a in managed memory explicitly, the OpenMP example also
>>>>> should work.
>>>>>
>>>>> There are other USM modes though, where the memory is shared implicitly
>>>>> between the host and the devices. Looks like currently LLVM
>>>>> implementation relies on this thing
>>>>>
>>>> https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#um-system-allocator
>>>>> where Implicit USM is supported.
>>>>>
>>>>> -------------
>>>>> Best regards,
>>>>> Alexey Bataev
>>>>>
>>>>> 3/1/2021 1:04 PM, Joachim Protze пишет:
>>>>>> Are the Kernel/Hardware requirements llvm specific?
>>>>>>
>>>>>> I can compile and execute the add_grid.cu example sucessfully:
>>>>>> https://developer.nvidia.com/blog/unified-memory-cuda-beginners/
>>>>>>
>>>>>> So, I would expect that an OpenMP program should also run sucessfully.
>>>>>>
>>>>>> - Joachim
>>>>>>
>>>>>>
>>>>>> Am 01.03.21 um 18:49 schrieb Alexey.Bataev:
>>>>>>> Hi, I you sure that you system supports Unified Shared Memory? As
>>>> far as
>>>>>>> I know it requires special linux kernel and the hardware must support
>>>>>>> it. If you system does not support it, the code will crash for sure
>>>> at
>>>>>>> the runtime.
>>>>>>>
>>>>>>> In this mode, IIRC, we just ignore map clauses since the accelerator
>>>>>>> devices can access the host memory directly without the need for
>>>>>>> allocating the device-specific memory.
>>>>>>>
>>>>>>>
>>>>>>> -------------
>>>>>>> Best regards,
>>>>>>> Alexey Bataev
>>>>>>>
>>>>>>> 3/1/2021 12:41 PM, Joachim Protze пишет:
>>>>>>>> Hi all,
>>>>>>>>
>>>>>>>> even a more simple example segfaults, when the requires directive
>>>> is there:
>>>>>>>>
>>>>>>>> #include <iostream>
>>>>>>>> #include <omp.h>
>>>>>>>> #include <stdio.h>
>>>>>>>>
>>>>>>>> #pragma omp requires unified_shared_memory
>>>>>>>> #define N 1024
>>>>>>>>
>>>>>>>> int main() {
>>>>>>>> int a[N];
>>>>>>>> printf("a=%p\n", a);
>>>>>>>> #pragma omp target map(tofrom : a[0:N])
>>>>>>>> {
>>>>>>>> printf("a=%p\n", a);
>>>>>>>> for (int i = 0; i < 1024; i++) {
>>>>>>>> a[i] = 123;
>>>>>>>> }
>>>>>>>> }
>>>>>>>> printf("a[0]=%i, a[%i]=%i\n", a[0], N/2, a[N/2]);
>>>>>>>> }
>>>>>>>>
>>>>>>>> The code runs sucessfully when the requires directive is removed
>>>> because
>>>>>>>> the mapping of `a` is explicitly specified.
>>>>>>>>
>>>>>>>> For this code to run successfully, would it be necessary to
>>>> allocate `a`
>>>>>>>> specially as cuda managed memory? I don't see any special treatment
>>>> of
>>>>>>>> `a` in llvm ir. As I understand the OpenMP spec, the requires
>>>> directive
>>>>>>>> should lead to a compile error if clang fails to generate such code.
>>>>>>>>
>>>>>>>> The requires example from the OpenMP Examples also fails with the
>>>> same
>>>>>>>> runtime error:
>>>>>>>>
>>>>>>>>
>>>> https://github.com/OpenMP/Examples/blob/main/sources/Example_requires.1.cpp
>>>>>>>>
>>>>>>>> - Joachim
>>>>>>>>
>>>>>>>> Am 28.02.21 um 11:12 schrieb Itaru Kitayama via Openmp-dev:
>>>>>>>>> This is the code:
>>>>>>>>>
>>>>>>>>> #include <iostream>
>>>>>>>>> #include <omp.h>
>>>>>>>>>
>>>>>>>>> #pragma omp requires unified_shared_memory
>>>>>>>>> #define N 1024
>>>>>>>>>
>>>>>>>>> int main() {
>>>>>>>>> int a[N] = {0};
>>>>>>>>> int *device_data = new int[N];
>>>>>>>>> #pragma omp target map(tofrom : device_data[0:N])
>>>>>>>>> {
>>>>>>>>> device_data = &a[0];
>>>>>>>>> for (int i = 0; i < 1024; i++) {
>>>>>>>>> device_data[i] = 123;
>>>>>>>>> }
>>>>>>>>> }
>>>>>>>>> std::cout << a[0] << std::endl;
>>>>>>>>> }
>>>>>>>>>
>>>>>>>>> On Sun, Feb 28, 2021 at 1:34 PM Johannes Doerfert
>>>>>>>>> <johannesdoerfert at gmail.com> wrote:
>>>>>>>>>> You have an illegal memory access, some memory is not properly
>>>>>>>>>> mapped.
>>>>>>>>>>
>>>>>>>>>>
>>>>>>>>>> On 2/27/21 7:47 PM, Itaru Kitayama wrote:
>>>>>>>>>>> Removed the internal function, but I get:
>>>>>>>>>>>
>>>>>>>>>>> CUDA device 0 info: Device supports up to 65536 CUDA blocks and
>>>> 1024
>>>>>>>>>>> threads with a warp size of 32
>>>>>>>>>>> CUDA device 0 info: Launching kernel
>>>>>>>>>>> __omp_offloading_34_8009dd23_main_l12 with 1 blocks and 33
>>>> threads in
>>>>>>>>>>> Generic mode
>>>>>>>>>>> CUDA error: Error when synchronizing stream. stream =
>>>>>>>>>>> 0x0000000001d22ae0, async info ptr = 0x00007ffe73ea2728
>>>>>>>>>>> CUDA error: an illegal memory access was encountered
>>>>>>>>>>> Libomptarget error: Failed to synchronize device.
>>>>>>>>>>> Libomptarget error: Call to targetDataEnd failed, abort target.
>>>>>>>>>>> Libomptarget error: Failed to process data after launching the
>>>> kernel.
>>>>>>>>>>> Libomptarget error: Source location information not present.
>>>> Compile
>>>>>>>>>>> with -g or -gline-tables-only.
>>>>>>>>>>> Libomptarget fatal error 1: failure of target construct while
>>>>>>>>>>> offloading is mandatory
>>>>>>>>>>> /var/spool/parastation/jobs/8941317: line 23: 20812 Aborted
>>>>>>>>>>> (core dumped) ./a.out
>>>>>>>>>>>
>>>>>>>>>>> On Sun, Feb 28, 2021 at 10:35 AM Alexey Bataev <
>>>> a.bataev at hotmail.com> wrote:
>>>>>>>>>>>> Do not call __tgt_register_requires directly, this is the
>>>> internal function called by global constructor and its arg value depends on
>>>> #pragma omp requires. Use just this pragma.
>>>>>>>>>>>>
>>>>>>>>>>>> Best regards,
>>>>>>>>>>>> Alexey Bataev
>>>>>>>>>>>>
>>>>>>>>>>>>> 27 февр. 2021 г., в 20:28, Itaru Kitayama via Openmp-dev <
>>>> openmp-dev at lists.llvm.org> написал(а):
>>>>>>>>>>>>>
>>>>>>>>>>>>> I'm trying to build a test C++ code that uses part of
>>>>>>>>>>>>> unified_shared_memory/shared_update.c
>>>>>>>>>>>>>
>>>>>>>>>>>>>> On Sun, Feb 28, 2021 at 10:25 AM Johannes Doerfert
>>>>>>>>>>>>>> <johannesdoerfert at gmail.com> wrote:
>>>>>>>>>>>>>>
>>>>>>>>>>>>>> I don't see this test, nor do I understand what you are
>>>> trying to say.
>>>>>>>>>>>>>> Is the test failing? If so, which test is this?
>>>>>>>>>>>>>>
>>>>>>>>>>>>>> ~ Johannes
>>>>>>>>>>>>>>
>>>>>>>>>>>>>>
>>>>>>>>>>>>>>> On 2/27/21 7:17 PM, Itaru Kitayama via Openmp-dev wrote:
>>>>>>>>>>>>>>> The below C++ code builds, but the executable fails at
>>>> runtime.
>>>>>>>>>>>>>>> (It is taken from the C code under the libomptarget subdir's
>>>> test directory)
>>>>>>>>>>>>>>>
>>>>>>>>>>>>>>> #include <omp.h>
>>>>>>>>>>>>>>>
>>>>>>>>>>>>>>> #pragma omp requires unified_shared_memory
>>>>>>>>>>>>>>> #define N 1024
>>>>>>>>>>>>>>> extern "C" void __tgt_register_requires(int64_t);
>>>>>>>>>>>>>>>
>>>>>>>>>>>>>>> int main() {
>>>>>>>>>>>>>>>
>>>>>>>>>>>>>>> int a[N] = {0};
>>>>>>>>>>>>>>> int b[N] = {0};
>>>>>>>>>>>>>>> int *device_data;
>>>>>>>>>>>>>>> __tgt_register_requires(1);
>>>>>>>>>>>>>>> #pragma omp target map(tofrom : device_data)
>>>>>>>>>>>>>>> {
>>>>>>>>>>>>>>> device_data = &a[0];
>>>>>>>>>>>>>>> for (int i = 0; i < 1024; i++) {
>>>>>>>>>>>>>>> a[i] += 1;
>>>>>>>>>>>>>>> }
>>>>>>>>>>>>>>> }
>>>>>>>>>>>>>>> }
>>>>>>>>>>>>>>> _______________________________________________
>>>>>>>>>>>>>>> Openmp-dev mailing list
>>>>>>>>>>>>>>> Openmp-dev at lists.llvm.org
>>>>>>>>>>>>>>> https://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
>>>>>>>>> _______________________________________________
>>>>>>>>> Openmp-dev mailing list
>>>>>>>>> Openmp-dev at lists.llvm.org
>>>>>>>>> https://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
>>
>
>
> _______________________________________________
> Openmp-dev mailing list
> Openmp-dev at lists.llvm.org
> https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev
>
More information about the Openmp-dev
mailing list