Dear Maintainer,

I found out that in the current debian gcc-14-offload-nvptx and
gcc-13-offload-nvptx, if I compile a code that requires
unified_shared_memory and uses openmp to offload to gpu, the code is never
run on the gpu.  It does compile the offload code, but then it is never
executed on the gpu.  If I attempt to run the code with
OMP_TARGET_OFFLOAD=MANDATORY, it gives the error:

libgomp: OMP_TARGET_OFFLOAD is set to MANDATORY, but only the host device is available

which is incorrect.  If I comment out the line

#pragma omp requires unified_shared_memory

then the code _does_ run on the gpu.

What is weird is that exactly the same code, on exactly the same hardware,
used to work in the past.  I cannot pinpoint exactly what update of what
package caused this issue.  I do know that the same code worked on my laptop
and still does offload to the gpu on another system with an old centos
system and hand-compiled gcc 13 with offload to nvptx-none.  Did anything
change in the required system setup and I missed it?  What puzzles me is
that it did use to work on exactly the same laptop, with sid, just some time
ago (one year?). And since the code _does_ compile and _does_ run without
errors unless one explicitly uses OMP_TARGET_OFFLOAD=MANDATORY, the change
may have occurred some time in the past and I did not notice till now.

I paste here a simple hello world code that shows the issue

#include <stdio.h>
#include <math.h>
#ifdef _OPENMP
#include <omp.h>
#endif
#define NX 100000000

#pragma omp requires unified_shared_memory

int main(void)
{
  double vecA[NX],vecB[NX],vecC[NX];
  double r=0.2;

/* Initialization of vectors */
#pragma omp target teams distribute parallel for simd
  for (long i = 0; i < NX; i++) {
     vecA[i] = pow(r, i);
     vecB[i] = 1.0;
  }

/* dot product of two vectors */
#pragma omp target teams distribute parallel for simd
  for (long i = 0; i < NX; i++) {
     vecC[i] = vecA[i] * vecB[i];
  }

  double sum = 0.0;
  /* calculate the sum */
  #pragma omp target teams distribute parallel for simd reduction(+:sum)
  for (long i = 0; i < NX; i++) {
    sum += vecC[i];
  }
  printf("The sum is: %8.6f \n", sum);
  return 0;
}

which I can compile with e.g.

gcc -O3 -fopenmp -foffload=nvptx-none -foffload-options="-O3 -fopt-info -lm" -o test test.c -lm

which compiles flawlessly and, due to the -fopt-info option given to the
offload compiler (and only to the offload compiler), gives out the following
info:

test.c:25:14: optimized: loop unrolled 7 times
test.c:32:9: optimized: loop unrolled 7 times
test.c:19:14: optimized: loop unrolled 3 times
test.c:30:11: optimized: basic block part vectorized using 16 byte vectors
test.c:23:9: optimized: basic block part vectorized using 16 byte vectors
test.c:16:9: optimized: basic block part vectorized using 16 byte vectors

if I run it (after allowing for a large stack with ulimit -s unlimited)
with just

./test

I get:

The sum is: 1.250000

if I run it with

OMP_TARGET_OFFLOAD=MANDATORY ./test

and the #pragma omp requires unified_shared_memory line was uncommented, I get

libgomp: OMP_TARGET_OFFLOAD is set to MANDATORY, but only the host device is available

while if I comment out #pragma omp requires unified_shared_memory recompile
and rerun with

OMP_TARGET_OFFLOAD=MANDATORY ./test

it does run on the gpu and produce the expected result. I actually verified
with nvidia-smi that it is indeed running on the gpu.

Thanks in advance for any info on how to get it to run again on the gpu
also when
#pragma omp requires unified_shared_memory
is enabled.

Best regards,
Giacomo Mulas

The OpenMP spec states that 'available devices' must be
'accessible' and 'supported'. And the later is defined
(glossary, here from 6.0):

"supported device - The host device or any non-host device supported
by the implementation, including any device-related requirements
specified by the requires directive."


Thus, if you specify

   omp requires unified_shared_memory

and either the device or the implementation does not support
unified-shared memory, all unsupported devices are removed such
that only the host is left (host fallback).

In some old GCC versions, '#pragma omp requires' was simply
ignored (warning with -Wunknown-pragmas, implied by -Wall).

For some versions, requiring USM would give an error.

I think since GCC 13, the host-fallback mechanism is at works,
printing an warning with GOMP_DEBUG=1 at runtime, if a device
cannot fulfill the requirement.

* * *

Since GCC 15, USM is supported under the following conditions:

https://gcc.gnu.org/onlinedocs/libgomp/Offload-Target-Specifics.html

As this is about an Nvidia GPU:

"* OpenMP code that has a requires directive with self_maps or
   unified_shared_memory runs on nvptx devices if and only if all
   of those support the pageableMemoryAccess property;⁵ otherwise,
   all nvptx device are removed from the list of available devices
   (“host fallback”)."

(5) https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#um-requirements


Which is fulfilled by [cf. (5)]:

"Linux HMM requires Linux kernel version 6.1.24+, 6.2.11+ or 6.3+,
devices with compute capability 7.5 or higher and a CUDA driver version
535+ installed with Open Kernel Modules." The pageableMemoryAccess is true on, e.g., the Frontier supercomputer
but also on my Laptop (compute capability 8.6, Ampere, meanwhile a
6.15 kernel), but admittedly we had some issues with Debian 12 Bookworm
and an Ada (8.9) card with the current 6.1.140 kernel (>= 6.1.24+) and
a recent open-kernel driver, even though it should have worked according
to the spec.

You can check this by something like:

   CUresult res;
   int n;
   res = cuInit (0);
   res = cuDeviceGetCount (&n);
   for (int dev = 0; dev < n; ++dev)
     {
       int val;
       __builtin_printf("============== Device %d =================\n", dev);
       res = cuDeviceGetAttribute (&val, CU_DEVICE_ATTRIBUTE_PAGEABLE_MEMORY_ACCESS, dev);
       __builtin_printf ("Device %d: pageableMemoryAccess: %d\n", dev, val);
     }

* * *

SOLUTION:

* For full USM support, your system needs to support
   pageableMemoryAccess (at least effectively for the devices involved)


If it does, you have two choices:

* Using GCC 15 from experimental, which supports USM, cf.
https://gcc.gnu.org/gcc-15/changes.html and
https://gcc.gnu.org/projects/gomp/#omp5.0 and
https://tracker.debian.org/pkg/gcc-15

* (Using (any) older GCC but) avoid using
     omp requires unified_shared_memory.

The difference between the two solutions:

- With the requirement, all maps are 'self maps', i.e.
   not data is actually copied.

- Without the requirement, data is copied but as, e.g.,
   pointer members of structs still point to the host memory,
   accessing those will work.

The USM support (HMM) works as follows:
If you access memory on the GPU that is not directly accessible
(= most host memory, unless you have e.g. a Grace-Hopper), a
memory-page fault is triggered and the the Linux kernel (+ Nvidia
kernel drivers) moves the page to device accessible memory.
Likewise on the way back from the device to the host accessible
memory.

* * *

If the system does not support pageableMemoryAccess, but at least
managedMemory you can access such memory (only) from the device.
If you are careful, this will work - but, obviously, the compiler
cannot regard such a system as supporting USM. Obtaining such memory
can be done using the CUDA-runtime routines for pinned and managed
memory.

I hope it helps.

Tobias

PS: GCC 16 will support some more memory handling tweaks and other
improvements. Still out to date and very much work in progress:
https://gcc.gnu.org/gcc-16/changes.html