[patch] OpenMP (C/C++): Keep pointer value of unmapped ptr with default mapping [PR110270]

[patch] OpenMP (C/C++): Keep pointer value of unmapped ptr with default mapping [PR110270]

[Tobias Burnus]

[-- Attachment #1: Type: text/plain, Size: 1855 bytes --]

This fixes an issue related to OpenMP C/C++'s default mapping of pointer variables.
(That's 'defaultmap(default:pointer)' – which is possibly surprisingly *not* the
same as 'defaultmap(firstprivate:pointer)').

Namely, OpenMP supports the following:

int *ptr = malloc(sizeof(int)*5);
#pragma omp target enter data map(ptr[:5])

#pragma omp target
   p[2] = 5;

which matches 'firstprivate(p)' + attaching the device address of 'p[:0]' (a zero-sized array),
the latter making it possible to use 'p' automatically without the need to add any map clauses
at least as long as *p has been mapped before.

However, for

  int *ptr = omp_target_alloc (sizeof(int)*5, dev_num);
  #pragma omp target
    p[2] = 5;
or for
   #pragma omp requires unified_shared_memory
   int pa = &A[0];
   #pragma omp target
     pa[0] = 6;

it failed before because neither 'ptr' nor 'pa' were mapped. Solution as a user was
either to add a (default)map clause (with map type than 'default'), a firstprivate
or an is_device_ptr clause.
The problem was that with default mapping, p and pa had the value NULL in the example
above on the device. (As required by OpenMP 5.0/5.1). With the commit, they retain
the original value avoiding surprises for the code above.
(See PR for the reference to the relevant sections of the OpenMP 5.{0,1,2} specifications.)


I would love if someone would give a review it;  albeit the actual code change in
libgomp/target.c is just a changing a single enum value.

If there are no comments, I intent to push it next week ...

Tobias
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

[-- Attachment #2: omp-ptr-priv.diff --]
[-- Type: text/x-patch, Size: 20020 bytes --]

OpenMP (C/C++): Keep pointer value of unmapped ptr with default mapping [PR110270]

For C/C++ pointers, default implicit mapping firstprivatizes the pointer
but if the memory it points to is mapped, the it is updated to point to
the device memory (by attaching a zero sized array section of the pointed-to
storage).

However, if the pointed-to storage wasn't mapped, the pointer was set to
NULL on the device side (OpenMP 5.0/5.1 semantic). With this commit, the
pointer retains the on-host address in that case (OpenMP 5.2 semantic).

The new semantic avoids an explicit map/firstprivate/is_device_ptr in the
following sensible cases: Special values (e.g. pointer or 0x1, 0x2 etc.),
explicitly device allocated memory (e.g. omp_target_alloc), and with
(unified) shared memory.
(Note: With (U)SM, mappings still must be tracked, at least when
omp_target_associate_ptr does not fail when passing in two destinct pointers.)

libgomp/

	PR middle-end/110270
	* target.c (gomp_map_vars_internal): Copy host value instead of NULL
	for  GOMP_MAP_ZERO_LEN_ARRAY_SECTION if not mapped.
	* libgomp.texi (OpenMP 5.2 Impl.): Mark as 'Y'.
	* testsuite/libgomp.c/target-19.c: Update expected value.
	* testsuite/libgomp.c++/target-18.C: Likewise.
	* testsuite/libgomp.c++/target-19.C: Likewise.
	* testsuite/libgomp.c-c++-common/requires-unified-addr-2.c: New test.
	* testsuite/libgomp.c-c++-common/target-implicit-map-3.c: New test.
	* testsuite/libgomp.c-c++-common/target-implicit-map-4.c: New test.

 libgomp/libgomp.texi                               |   2 +-
 libgomp/target.c                                   |   2 +-
 libgomp/testsuite/libgomp.c++/target-18.C          |  21 ++-
 libgomp/testsuite/libgomp.c++/target-19.C          |  13 +-
 .../libgomp.c-c++-common/requires-unified-addr-2.c |  85 +++++++++++
 .../libgomp.c-c++-common/target-implicit-map-3.c   | 105 ++++++++++++++
 .../libgomp.c-c++-common/target-implicit-map-4.c   | 159 +++++++++++++++++++++
 libgomp/testsuite/libgomp.c/target-19.c            |  21 ++-
 8 files changed, 392 insertions(+), 16 deletions(-)

diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 1c57f5aa261..db8b1f1427e 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -384,7 +384,7 @@ to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab
 @item @code{declare mapper} with iterator and @code{present} modifiers
       @tab N @tab
 @item If a matching mapped list item is not found in the data environment, the
-      pointer retains its original value @tab N @tab
+      pointer retains its original value @tab Y @tab
 @item New @code{enter} clause as alias for @code{to} on declare target directive
       @tab Y @tab
 @item Deprecation of @code{to} clause on declare target directive @tab N @tab
diff --git a/libgomp/target.c b/libgomp/target.c
index e39ef8f6e82..aa2410c0f16 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1149,7 +1149,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 	  if (!n)
 	    {
 	      tgt->list[i].key = NULL;
-	      tgt->list[i].offset = OFFSET_POINTER;
+	      tgt->list[i].offset = OFFSET_INLINED;
 	      continue;
 	    }
 	}
diff --git a/libgomp/testsuite/libgomp.c++/target-18.C b/libgomp/testsuite/libgomp.c++/target-18.C
index f1085b14022..a21ed4e81f9 100644
--- a/libgomp/testsuite/libgomp.c++/target-18.C
+++ b/libgomp/testsuite/libgomp.c++/target-18.C
@@ -20,7 +20,9 @@ foo (int *&p, int *&q, int *&r, int n, int m)
 	  err = 1;
       if (sep)
 	{
-	  if (q != (int *) 0 || r != (int *) 0)
+	  /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to
+	     unmapped storage.  */
+	  if (q == (int *) 0 || r == (int *) 0)
 	    err = 1;
 	}
       else if (p + 8 != q || r != s)
@@ -37,7 +39,9 @@ foo (int *&p, int *&q, int *&r, int n, int m)
 	  err = 1;
       if (sep)
 	{
-	  if (q != (int *) 0 || r != (int *) 0)
+	  /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to
+	     unmapped storage.  */
+	  if (q == (int *) 0 || r == (int *) 0)
 	    err = 1;
 	}
       else if (p + 8 != q || r != s)
@@ -55,7 +59,9 @@ foo (int *&p, int *&q, int *&r, int n, int m)
 	  err = 1;
       if (sep)
 	{
-	  if (q != (int *) 0 || r != (int *) 0)
+	  /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to
+	     unmapped storage.  */
+	  if (q == (int *) 0 || r == (int *) 0)
 	    err = 1;
 	}
       else if (p + 8 != q || r != s)
@@ -91,7 +97,8 @@ foo (int *&p, int *&q, int *&r, int n, int m)
 	  err = 1;
 	else if (sep)
 	  {
-	    if (r != (int *) 0)
+	    /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/
+	    if (r == (int *) 0)
 	      err = 1;
 	  }
 	else if (r != q + 1)
@@ -110,7 +117,8 @@ foo (int *&p, int *&q, int *&r, int n, int m)
 	  err = 1;
 	else if (sep)
 	  {
-	    if (r != (int *) 0)
+	    /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/
+	    if (r == (int *) 0)
 	      err = 1;
 	  }
 	else if (r != q + 1)
@@ -130,7 +138,8 @@ foo (int *&p, int *&q, int *&r, int n, int m)
 	  err = 1;
 	else if (sep)
 	  {
-	    if (r != (int *) 0)
+	    /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/
+	    if (r == (int *) 0)
 	      err = 1;
 	  }
 	else if (r != q + 1)
diff --git a/libgomp/testsuite/libgomp.c++/target-19.C b/libgomp/testsuite/libgomp.c++/target-19.C
index afa6e68d5cc..7bae31d2734 100644
--- a/libgomp/testsuite/libgomp.c++/target-19.C
+++ b/libgomp/testsuite/libgomp.c++/target-19.C
@@ -1,3 +1,8 @@
+/* { dg-additional-options "-O0" } */
+/* Disable optimization to ensure that the compiler does not exploit that
+   S::r + t will never be NULL due to int (&r) and (&t).  */
+
+
 extern "C" void abort ();
 struct S { char a[64]; int (&r)[2]; char b[64]; };
 
@@ -19,7 +24,9 @@ foo (S s, int (&t)[3], int z)
   #pragma omp target map(from: err) map(tofrom: s.r[:0], t[:0])
   {
     if (sep)
-      err = s.r != (int *) 0 || t != (int *) 0;
+      /* Since OpenMP 5.2, if no matching mapped list it has been found,
+	 pointers retain their original value.  */
+      err = s.r == (int *) 0 || t == (int *) 0;
     else
       err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7;
   }
@@ -28,7 +35,9 @@ foo (S s, int (&t)[3], int z)
   #pragma omp target map(from: err) map(tofrom: s.r[:z], t[:z])
   {
     if (sep)
-      err = s.r != (int *) 0 || t != (int *) 0;
+      /* Since OpenMP 5.2, if no matching mapped list it has been found,
+	 pointers retain their original value.  */
+      err = s.r == (int *) 0 || t == (int *) 0;
     else
       err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7;
   }
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-2.c b/libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-2.c
new file mode 100644
index 00000000000..3b5dcd38c1a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-2.c
@@ -0,0 +1,85 @@
+/* PR middle-end/110270 */
+
+/* OpenMP 5.2's 'defaultmap(default : pointer) for C/C++ pointers retains the
+   pointer value instead of setting it to NULL if the pointer cannot be found.
+   Contrary to requires-unified-addr-1.c which is valid OpenMP 5.0/5.1/5.2,
+   this testcase is only valid since OpenMP 5.2.  */
+
+/* This is kind of a follow-up to the requires-unified-addr-1.c testcase
+   and  PR libgomp/109837 */
+
+
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#pragma omp requires unified_address
+
+#define N 15
+
+void
+test_device (int dev)
+{
+  struct st {
+    int *ptr;
+    int n;
+  };
+  struct st s;
+
+  s.n = 10;
+  s.ptr = (int *) omp_target_alloc (sizeof (int)*s.n, dev);
+  int *ptr1 = (int *) omp_target_alloc (sizeof (int)*N, dev);
+  assert (s.ptr != NULL);
+  assert (ptr1 != NULL);
+
+  int q[4] = {1,2,3,4};
+  int *qptr;
+  #pragma omp target enter data map(q) device(device_num: dev)
+  #pragma omp target data use_device_addr(q) device(device_num: dev)
+    qptr = q;
+
+  #pragma omp target map(to:s) device(device_num: dev)
+  for (int i = 0; i < s.n; i++)
+    s.ptr[i] = 23*i;
+
+  int *ptr2 = &s.ptr[3];
+
+  /* s.ptr is not mapped (but omp_target_alloc'ed) thus ptr2 shall retain its value.  */
+  #pragma omp target device(device_num: dev)  /* implied: defaultmap(default : pointer) */
+  for (int i = 0; i < 4; i++)
+    *(qptr++) = ptr2[i];
+
+  #pragma omp target exit data map(q) device(device_num: dev)
+  for (int i = 0; i < 4; i++)
+    q[i] = 23 * (i+3);
+
+  /* ptr1 retains the value as it is not mapped (but it is omp_target_alloc'ed). */
+  #pragma omp target defaultmap(default : pointer) device(device_num: dev)
+  for (int i = 0; i < N; i++)
+    ptr1[i] = 11*i;
+
+  int *ptr3 = (int *) malloc (sizeof (int)*N);
+  assert (0 == omp_target_memcpy(ptr3, ptr1, N * sizeof(int), 0, 0,
+				 omp_get_initial_device(), dev));
+  for (int i = 0; i < N; i++)
+    assert (ptr3[i] == 11*i);
+
+  free (ptr3);
+  omp_target_free (ptr1, dev);
+  omp_target_free (s.ptr, dev);
+}
+
+int
+main()
+{
+  int ntgts = omp_get_num_devices();
+  if (ntgts)
+    fprintf (stderr, "Offloading devices exist\n");  /* { dg-output "Offloading devices exist(\n|\r\n|\r)" { target offload_device } } */
+  else
+    fprintf (stderr, "Only host fallback\n");        /* { dg-output "Only host fallback(\n|\r\n|\r)" { target { ! offload_device } } } */
+
+  for (int i = 0; i <= ntgts; i++)
+    test_device (i);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c
new file mode 100644
index 00000000000..e69f6d0db2f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c
@@ -0,0 +1,105 @@
+/* PR middle-end/110270 */
+
+/* Ensure that defaultmap(default : pointer) uses correct OpenMP 5.2
+   semantic, i.e. keeping the pointer value even if not mapped;
+   before OpenMP 5.0/5.1 required that it is NULL, causing issues
+   especially with unified-shared memory but also the code below
+   shows why that's not a good idea.  */
+
+#include <stdio.h>
+#include <stdint.h>
+#include <omp.h>
+
+/* 'unified_address' is required by the OpenMP spec as only then
+   'is_device_ptr' can be left out.  All our devices support this
+   while remote offloading would not.  However, in practice it is
+   sufficient that the host and device pointer size is the same
+   (or the device pointer is smaller) - and then a device pointer is
+   representable and omp_target_alloc can return a bare device pointer.
+
+   We here assume that this weaker condition holds and do not
+   require:   #pragma omp requires unified_address  */
+
+void
+test_device (int dev)
+{
+  int *p1 = (int*) 0x12345;
+  int *p1a = (int*) 0x67890;
+  int *p2 = (int*) omp_target_alloc (sizeof (int) * 5, dev);
+  int *p2a = (int*) omp_target_alloc (sizeof (int) * 10, dev);
+  intptr_t ip = (intptr_t) p2;
+  intptr_t ipa = (intptr_t) p2a;
+
+  int A[3] = {1,2,3};
+  int B[5] = {4,5,6,7,8};
+  int *p3 = &A[0];
+  int *p3a = &B[0];
+
+  #pragma omp target enter data map(to:A) device(dev)
+
+  #pragma omp target device(dev) /* defaultmap(default:pointer) */
+  {
+    /* The pointees aren't mapped. */
+    /* OpenMP 5.2 -> same value as before the target region. */
+    if ((intptr_t) p1 != 0x12345) __builtin_abort ();
+    if ((intptr_t) p2 != ip) __builtin_abort ();  
+    for (int i = 0; i < 5; i++)
+      p2[i] = 13*i;
+
+    for (int i = 0; i < 10; i++)
+      ((int *)ipa)[i] = 7*i;
+
+    /* OpenMP: Mapped => must point to the corresponding device storage of 'A' */
+    if (p3[0] != 1 || p3[1] != 2 || p3[2] != 3)
+      __builtin_abort ();
+    p3[0] = -11; p3[1] = -22; p3[2] = -33;
+  }
+  #pragma omp target exit data map(from:A) device(dev)
+
+  if (p3[0] != -11 || p3[1] != -22 || p3[2] != -33)
+    __builtin_abort ();
+
+  // With defaultmap:
+
+  #pragma omp target enter data map(to:B) device(dev)
+
+  #pragma omp target device(dev) defaultmap(default:pointer)
+  {
+    /* The pointees aren't mapped. */
+    /* OpenMP 5.2 -> same value as before the target region. */
+    if ((intptr_t) p1a != 0x67890) __builtin_abort ();
+    if ((intptr_t) p2a != ipa) __builtin_abort ();  
+
+    for (int i = 0; i < 5; i++)
+      ((int *)ip)[i] = 13*i;
+
+    for (int i = 0; i < 10; i++)
+      p2a[i] = 7*i;
+
+    /* OpenMP: Mapped => must point to the corresponding device storage of 'B' */
+    if (p3a[0] != 4 || p3a[1] != 5 || p3a[2] != 6 || p3a[3] != 7 || p3a[4] != 8)
+      __builtin_abort ();
+    p3a[0] = -44; p3a[1] = -55; p3a[2] = -66; p3a[3] = -77; p3a[4] = -88;
+  }
+  #pragma omp target exit data map(from:B) device(dev)
+
+  if (p3a[0] != -44 || p3a[1] != -55 || p3a[2] != -66 || p3a[3] != -77 || p3a[4] != -88)
+    __builtin_abort ();
+
+  omp_target_free (p2, dev);
+  omp_target_free (p2a, dev);
+}
+
+int
+main()
+{
+  int ntgts = omp_get_num_devices();
+  if (ntgts)
+    fprintf (stderr, "Offloading devices exist\n");  /* { dg-output "Offloading devices exist(\n|\r\n|\r)" { target offload_device } } */
+  else
+    fprintf (stderr, "Only host fallback\n");        /* { dg-output "Only host fallback(\n|\r\n|\r)" { target { ! offload_device } } } */
+
+  for (int i = 0; i <= ntgts; i++)
+    test_device (i);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c
new file mode 100644
index 00000000000..f92abfdb30d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c
@@ -0,0 +1,159 @@
+/* PR middle-end/110270 */
+
+/* Same as target-implicit-map-3.c but uses the following requiement
+   and for not mapping the stack variables 'A' and 'B' (not mapped
+   but accessible -> USM makes this tested feature even more important.)  */
+
+#pragma omp requires unified_shared_memory
+
+/* Ensure that defaultmap(default : pointer) uses correct OpenMP 5.2
+   semantic, i.e. keeping the pointer value even if not mapped;
+   before OpenMP 5.0/5.1 required that it is NULL. */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdint.h>
+#include <omp.h>
+
+void
+test_device (int dev)
+{
+  int *p1 = (int*) 0x12345;
+  int *p1a = (int*) 0x67890;
+  int *p2 = (int*) omp_target_alloc (sizeof (int) * 5, dev);
+  int *p2a = (int*) omp_target_alloc (sizeof (int) * 10, dev);
+  intptr_t ip = (intptr_t) p2;
+  intptr_t ipa = (intptr_t) p2a;
+
+  int A[3] = {1,2,3};
+  int B[5] = {4,5,6,7,8};
+  int *p3 = &A[0];
+  int *p3a = &B[0];
+
+  const omp_alloctrait_t traits[]
+    = { { omp_atk_alignment, 128 },
+        { omp_atk_pool_size, 1024 }};
+  omp_allocator_handle_t a = omp_init_allocator (omp_default_mem_space, 2, traits);
+
+  int *p4 = (int*) malloc (sizeof (int) * 5);
+  int *p4a = (int*) omp_alloc (sizeof (int) * 10, a);
+  intptr_t ip4 = (intptr_t) p4;
+  intptr_t ip4a = (intptr_t) p4a;
+
+  for (int i = 0; i < 5; i++)
+    p4[i] = -31*i;
+
+  for (int i = 0; i < 10; i++)
+    p4a[i] = -43*i;
+
+  /* Note: 'A' is not mapped but USM accessible.  */
+  #pragma omp target device(dev) /* defaultmap(default:pointer) */
+  {
+    /* The pointees aren't mapped. */
+    /* OpenMP 5.2 -> same value as before the target region. */
+    if ((intptr_t) p1 != 0x12345) abort ();
+    if ((intptr_t) p2 != ip) abort ();  
+    for (int i = 0; i < 5; i++)
+      p2[i] = 13*i;
+
+    for (int i = 0; i < 10; i++)
+      ((int *)ipa)[i] = 7*i;
+
+    /* OpenMP: Points to 'A'. */
+    if (p3[0] != 1 || p3[1] != 2 || p3[2] != 3)
+      abort ();
+    p3[0] = -11; p3[1] = -22; p3[2] = -33;
+
+    /* USM accesible allocated host memory. */
+    if ((intptr_t) p4 != ip4)
+      abort ();  
+    for (int i = 0; i < 5; i++)
+      if (p4[i] != -31*i)
+	abort ();
+    for (int i = 0; i < 10; i++)
+      if (((int *)ip4a)[i] != -43*i)
+	abort ();
+    for (int i = 0; i < 5; i++)
+      p4[i] = 9*i;
+    for (int i = 0; i < 10; i++)
+      ((int *)ip4a)[i] = 18*i;
+  }
+
+  if (p3[0] != -11 || p3[1] != -22 || p3[2] != -33)
+    abort ();
+
+  for (int i = 0; i < 5; i++)
+    if (p4[i] != 9*i)
+      abort ();
+  for (int i = 0; i < 10; i++)
+    if (p4a[i] != 18*i)
+      abort ();
+  for (int i = 0; i < 5; i++)
+    p4[i] = -77*i;
+  for (int i = 0; i < 10; i++)
+    p4a[i] = -65*i;
+
+  // With defaultmap:
+
+  /* Note: 'B' is not mapped but USM accessible.  */
+  #pragma omp target device(dev) defaultmap(default:pointer)
+  {
+    /* The pointees aren't mapped. */
+    /* OpenMP 5.2 -> same value as before the target region. */
+    if ((intptr_t) p1a != 0x67890) abort ();
+    if ((intptr_t) p2a != ipa) abort ();  
+
+    for (int i = 0; i < 5; i++)
+      ((int *)ip)[i] = 13*i;
+
+    for (int i = 0; i < 10; i++)
+      p2a[i] = 7*i;
+
+    /* USM accesible allocated host memory. */
+    if ((intptr_t) p4a != ip4a) abort ();  
+
+    /* OpenMP: Points to 'B'. */
+    if (p3a[0] != 4 || p3a[1] != 5 || p3a[2] != 6 || p3a[3] != 7 || p3a[4] != 8)
+      abort ();
+    p3a[0] = -44; p3a[1] = -55; p3a[2] = -66; p3a[3] = -77; p3a[4] = -88;
+
+    /* USM accesible allocated host memory. */
+    if ((intptr_t) p4a != ip4a)
+      abort ();  
+    for (int i = 0; i < 5; i++)
+      if (((int *)ip4)[i] != -77*i)
+	abort ();
+    for (int i = 0; i < 10; i++)
+      if (p4a[i] != -65*i)
+	abort ();
+    for (int i = 0; i < 5; i++)
+      p4[i] = 36*i;
+    for (int i = 0; i < 10; i++)
+      ((int *)ip4a)[i] = 4*i;
+  }
+
+  if (p3a[0] != -44 || p3a[1] != -55 || p3a[2] != -66 || p3a[3] != -77 || p3a[4] != -88)
+    abort ();
+
+  for (int i = 0; i < 5; i++)
+    if (p4[i] != 36*i)
+      abort ();
+  for (int i = 0; i < 10; i++)
+    if (p4a[i] != 4*i)
+      abort ();
+
+  omp_target_free (p2, dev);
+  omp_target_free (p2a, dev);
+  free (p4);
+  omp_free (p4a, a);
+  omp_destroy_allocator (a);
+}
+
+int
+main()
+{
+  int ntgts = omp_get_num_devices();
+  for (int i = 0; i <= ntgts; i++)
+    test_device (i);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-19.c b/libgomp/testsuite/libgomp.c/target-19.c
index 2505cafca9f..dac7c56cbde 100644
--- a/libgomp/testsuite/libgomp.c/target-19.c
+++ b/libgomp/testsuite/libgomp.c/target-19.c
@@ -20,7 +20,9 @@ foo (int *p, int *q, int *r, int n, int m)
 	  err = 1;
       if (sep)
 	{
-	  if (q != (int *) 0 || r != (int *) 0)
+	  /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to
+	     unmapped storage.  */
+	  if (q == (int *) 0 || r == (int *) 0)
 	    err = 1;
 	}
       else if (p + 8 != q || r != s)
@@ -37,7 +39,9 @@ foo (int *p, int *q, int *r, int n, int m)
 	  err = 1;
       if (sep)
 	{
-	  if (q != (int *) 0 || r != (int *) 0)
+	  /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to
+	     unmapped storage.  */
+	  if (q == (int *) 0 || r == (int *) 0)
 	    err = 1;
 	}
       else if (p + 8 != q || r != s)
@@ -55,7 +59,9 @@ foo (int *p, int *q, int *r, int n, int m)
 	  err = 1;
       if (sep)
 	{
-	  if (q != (int *) 0 || r != (int *) 0)
+	  /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to
+	     unmapped storage.  */
+	  if (q == (int *) 0 || r == (int *) 0)
 	    err = 1;
 	}
       else if (p + 8 != q || r != s)
@@ -91,7 +97,8 @@ foo (int *p, int *q, int *r, int n, int m)
 	  err = 1;
 	else if (sep)
 	  {
-	    if (r != (int *) 0)
+	    /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/
+	    if (r == (int *) 0)
 	      err = 1;
 	  }
 	else if (r != q + 1)
@@ -110,7 +117,8 @@ foo (int *p, int *q, int *r, int n, int m)
 	  err = 1;
 	else if (sep)
 	  {
-	    if (r != (int *) 0)
+	    /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/
+	    if (r == (int *) 0)
 	      err = 1;
 	  }
 	else if (r != q + 1)
@@ -130,7 +138,8 @@ foo (int *p, int *q, int *r, int n, int m)
 	  err = 1;
 	else if (sep)
 	  {
-	    if (r != (int *) 0)
+	    /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/
+	    if (r == (int *) 0)
 	      err = 1;
 	  }
 	else if (r != q + 1)