[Openmp-dev] Libomptarget fatal error 1: '#pragma omp requires unified_shared_memory' not used consistently!

[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