OpenMP: Fix mapping of zero-sized arrays with non-literal size: map(var[:n]), n = 0devel/omp/gcc-14

For map(ptr[:0]), the used map kind is GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION
and it is permitted that 'ptr' does not exist. 'ptr' is set to the device
pointee if it exists or to the host value otherwise.

For map(ptr[:3]), the variable is first mapped and then ptr is updated to point
to the just-mapped device data; the attachment uses GOMP_MAP_ATTACH.

For map(ptr[:n]), generates always a GOMP_MAP_ATTACH, but when n == 0, it
was failing with:
   "pointer target not mapped for attach"

The solution is not to fail but first to check whether it was mapped before.
It turned out that for the mapping part, GCC adds a run-time check whether
n == 0 - and uses GOMP_MAP_ZERO_LEN_ARRAY_SECTION for the mapping.
Thus, we just have to check whether there such a mapping for the address
for which the GOMP_MAP_ATTACH. was requested. And, if there was, the
error diagnostic can be skipped.

Unsurprisingly, this issue occurs in real-world code; it was detected in
a code that distributes work via MPI and for some processes, some bounds
ended up to be zero.

libgomp/ChangeLog:

	* target.c (gomp_attach_pointer): Return bool; accept additional
	bool to optionally silence the fatal pointee-not-found error.
	(gomp_map_vars_internal): If the pointee could not be found,
	check whether it was mapped as GOMP_MAP_ZERO_LEN_ARRAY_SECTION.
	* libgomp.h (gomp_attach_pointer): Update prototype.
	* oacc-mem.c (acc_attach_async, goacc_enter_data_internal): Update
	calls.
	* testsuite/libgomp.c/target-map-zero-sized.c: New test.
	* testsuite/libgomp.c/target-map-zero-sized-2.c: New test.
	* testsuite/libgomp.c/target-map-zero-sized-3.c: New test.

(cherry picked from commit 814e29e390b1e9253f9a38e0d84f5ebe5de0c13e)

7 files changed, 304 insertions, 16 deletions

diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 264a4d7..a6d676c 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,19 @@
+2025-05-14  Tobias Burnus  <tburnus@baylibre.com>
+
+	Backported from master:
+	2025-05-14  Tobias Burnus  <tburnus@baylibre.com>
+
+	* target.c (gomp_attach_pointer): Return bool; accept additional
+	bool to optionally silence the fatal pointee-not-found error.
+	(gomp_map_vars_internal): If the pointee could not be found,
+	check whether it was mapped as GOMP_MAP_ZERO_LEN_ARRAY_SECTION.
+	* libgomp.h (gomp_attach_pointer): Update prototype.
+	* oacc-mem.c (acc_attach_async, goacc_enter_data_internal): Update
+	calls.
+	* testsuite/libgomp.c/target-map-zero-sized.c: New test.
+	* testsuite/libgomp.c/target-map-zero-sized-2.c: New test.
+	* testsuite/libgomp.c/target-map-zero-sized-3.c: New test.
+
 2025-04-25  Thomas Schwinge  <tschwinge@baylibre.com>
 
 	Backported from trunk:
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 43e7a06..a60a3d8 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1487,10 +1487,10 @@ extern void gomp_copy_dev2host (struct gomp_device_descr *,
 				struct goacc_asyncqueue *, void *, const void *,
 				size_t);
 extern uintptr_t gomp_map_val (struct target_mem_desc *, void **, size_t);
-extern void gomp_attach_pointer (struct gomp_device_descr *,
+extern bool gomp_attach_pointer (struct gomp_device_descr *,
 				 struct goacc_asyncqueue *, splay_tree,
 				 splay_tree_key, uintptr_t, size_t,
-				 struct gomp_coalesce_buf *, bool);
+				 struct gomp_coalesce_buf *, bool, bool);
 extern void gomp_detach_pointer (struct gomp_device_descr *,
 				 struct goacc_asyncqueue *, splay_tree_key,
 				 uintptr_t, bool, struct gomp_coalesce_buf *);
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index feca3b2..a4ddb7d 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -979,7 +979,7 @@ acc_attach_async (void **hostaddr, int async)
     }
 
   gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, (uintptr_t) hostaddr,
-		       0, NULL, false);
+		       0, NULL, false, true);
 
   gomp_mutex_unlock (&acc_dev->lock);
 }
@@ -1215,7 +1215,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 	  if ((kinds[i] & 0xff) == GOMP_MAP_ATTACH)
 	    {
 	      gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n,
-				   (uintptr_t) h, s, NULL, false);
+				   (uintptr_t) h, s, NULL, false, true);
 	      /* OpenACC 'attach'/'detach' doesn't affect structured/dynamic
 		 reference counts ('n->refcount', 'n->dynamic_refcount').  */
 	    }
@@ -1233,7 +1233,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 		  = lookup_host (acc_dev, hostaddrs[j], sizeof (void *));
 		gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, m,
 				     (uintptr_t) hostaddrs[j], sizes[j], NULL,
-				     false);
+				     false, true);
 	      }
 
 	  bool processed = false;
diff --git a/libgomp/target.c b/libgomp/target.c
index 796decd..ab1fae1 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -800,12 +800,22 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
 	      (void *) cur_node.host_end);
 }
 
-attribute_hidden void
+/* Update the devptr by setting it to the device address of the host pointee
+   'attach_to'; devptr is obtained from the splay_tree_key n.
+   When the pointer is already attached or the host pointee is either
+   NULL or in memory map, this function returns true.
+   Otherwise, the device pointer is set to point to the host pointee and:
+   - If allow_zero_length_array_sections is set, true is returned.
+   - Else, if fail_if_not_found is set, a fatal error is issued.
+   - Otherwise, false is returned.  */
+
+attribute_hidden bool
 gomp_attach_pointer (struct gomp_device_descr *devicep,
 		     struct goacc_asyncqueue *aq, splay_tree mem_map,
 		     splay_tree_key n, uintptr_t attach_to, size_t bias,
 		     struct gomp_coalesce_buf *cbufp,
-		     bool allow_zero_length_array_sections)
+		     bool allow_zero_length_array_sections,
+		     bool fail_if_not_found)
 {
   struct splay_tree_key_s s;
   size_t size, idx;
@@ -860,7 +870,7 @@ gomp_attach_pointer (struct gomp_device_descr *devicep,
 	  gomp_copy_host2dev (devicep, aq, (void *) devptr, (void *) &data,
 			      sizeof (void *), true, cbufp);
 
-	  return;
+	  return true;
 	}
 
       s.host_start = target + bias;
@@ -869,15 +879,16 @@ gomp_attach_pointer (struct gomp_device_descr *devicep,
 
       if (!tn)
 	{
-	  if (allow_zero_length_array_sections)
-	    /* When allowing attachment to zero-length array sections, we
-	       copy the host pointer when the target region is not mapped.  */
-	    data = target;
-	  else
+	  /* We copy the host pointer when the target region is not mapped;
+	     for allow_zero_length_array_sections, that's permitted.
+	     Otherwise, it depends on the context. Return false in that
+	     case, unless fail_if_not_found.  */
+	  if (!allow_zero_length_array_sections && fail_if_not_found)
 	    {
 	      gomp_mutex_unlock (&devicep->lock);
 	      gomp_fatal ("pointer target not mapped for attach");
 	    }
+	  data = target;
 	}
       else
 	data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start;
@@ -889,10 +900,13 @@ gomp_attach_pointer (struct gomp_device_descr *devicep,
 
       gomp_copy_host2dev (devicep, aq, (void *) devptr, (void *) &data,
 			  sizeof (void *), true, cbufp);
+      if (!tn && !allow_zero_length_array_sections)
+	return false;
     }
   else
     gomp_debug (1, "%s: attach count for %p -> %u\n", __FUNCTION__,
 		(void *) attach_to, (int) n->aux->attach_count[idx]);
+  return true;
 }
 
 attribute_hidden void
@@ -1794,9 +1808,37 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 		      bool zlas
 			= ((kind & typemask)
 			   == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
-		      gomp_attach_pointer (devicep, aq, mem_map, n,
-					   (uintptr_t) hostaddrs[i], sizes[i],
-					   cbufp, zlas);
+		      /* For 'target enter data', the map clauses are split;
+			 however, for more complex code with struct and
+			 pointer members, the mapping and the attach can end up
+			 in different sets; or the wrong mapping with the
+			 attach. As there is no way to know whether a size
+			 zero like  'var->ptr[i][:0]' happend in the same
+			 directive or not, the not-attached check is now
+			 fully silenced for 'enter data'.  */
+		      if (openmp_p && (pragma_kind & GOMP_MAP_VARS_ENTER_DATA))
+			zlas = true;
+		      if (!gomp_attach_pointer (devicep, aq, mem_map, n,
+						(uintptr_t) hostaddrs[i], sizes[i],
+						cbufp, zlas, !openmp_p))
+			{
+			  /* Pointee not found; that's an error except for
+			     map(var[:n]) with n == 0; the compiler adds a
+			     runtime condition such that for those the kind is
+			     always GOMP_MAP_ZERO_LEN_ARRAY_SECTION.  */
+			  for (j = i; j > 0; j--)
+			    if (*(void**) hostaddrs[i] == hostaddrs[j-1] - sizes[i]
+				&& sizes[j-1] == 0
+				&& (GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+				    == (get_kind (short_mapkind, kinds, j-1)
+					& typemask)))
+			      break;
+			  if (j == 0)
+			    {
+			      gomp_mutex_unlock (&devicep->lock);
+			      gomp_fatal ("pointer target not mapped for attach");
+			    }
+			}
 		    }
 		  else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0)
 		    {
diff --git a/libgomp/testsuite/libgomp.c/target-map-zero-sized-2.c b/libgomp/testsuite/libgomp.c/target-map-zero-sized-2.c
new file mode 100644
index 0000000..3220828
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-map-zero-sized-2.c
@@ -0,0 +1,74 @@
+int
+main ()
+{
+  int i, n;
+  int data[] = {1,2};
+  struct S { int **ptrset; };
+
+// -----------------------------------
+
+/* The produced mapping for sptr1->ptrset[i][:n]
+
+   GOMP_MAP_STRUCT (size = 1)
+      GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+   GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+      GOMP_MAP_ATTACH
+   GOMP_MAP_ATTACH -> attaching to 2nd GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+
+which get split into 3 separate map_vars call; in particular,
+the latter is separate and points to an unmpapped variable.
+
+Thus, it failed with:
+   libgomp: pointer target not mapped for attach  */
+
+  struct S s1, *sptr1;
+  s1.ptrset = (int **) __builtin_malloc (sizeof(void*) * 3);
+  s1.ptrset[0] = data;
+  s1.ptrset[1] = data;
+  s1.ptrset[2] = data;
+  sptr1 = &s1;
+
+  i = 1;
+  n = 0;
+  #pragma omp target enter data map(sptr1[:1], sptr1->ptrset[:3])
+  #pragma omp target enter data map(sptr1->ptrset[i][:n])
+
+  #pragma omp target exit data map(sptr1->ptrset[i][:n])
+  #pragma omp target exit data map(sptr1[:1], sptr1->ptrset[:3])
+
+  __builtin_free (s1.ptrset);
+
+// -----------------------------------
+
+/* The produced mapping for sptr2->ptrset[i][:n] is similar:
+
+   GOMP_MAP_STRUCT (size = 1)
+      GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+   GOMP_MAP_TO  ! this one has now a finite size
+      GOMP_MAP_ATTACH
+   GOMP_MAP_ATTACH -> attach to the GOMP_MAP_TO
+
+As the latter GOMP_MAP_ATTACH has now a pointer target,
+the attachment worked.  */
+
+  struct S s2, *sptr2;
+  s2.ptrset = (int **) __builtin_malloc (sizeof(void*) * 3);
+  s2.ptrset[0] = data;
+  s2.ptrset[1] = data;
+  s2.ptrset[2] = data;
+  sptr2 = &s2;
+
+  i = 1;
+  n = 2;
+  #pragma omp target enter data map(sptr2[:1], sptr2->ptrset[:3])
+  #pragma omp target enter data map(sptr2->ptrset[i][:n])
+
+  #pragma omp target
+    if (sptr2->ptrset[1][0] != 1 || sptr2->ptrset[1][1] != 2)
+      __builtin_abort ();
+
+  #pragma omp target exit data map(sptr2->ptrset[i][:n])
+  #pragma omp target exit data map(sptr2[:1], sptr2->ptrset[:3])
+
+  __builtin_free (s2.ptrset);
+}
diff --git a/libgomp/testsuite/libgomp.c/target-map-zero-sized-3.c b/libgomp/testsuite/libgomp.c/target-map-zero-sized-3.c
new file mode 100644
index 0000000..f968bd3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-map-zero-sized-3.c
@@ -0,0 +1,49 @@
+int
+main ()
+{
+  int i, n, n2;
+  int data[] = {1,2};
+  struct S {
+    int **ptrset;
+    int **ptrset2;
+  };
+
+  /* This is the same as target-map-zero-sized-3.c, but by mixing
+     mapped and non-mapped items, the mapping before the ATTACH
+     might (or here: is) not actually associated with the the
+     pointer used for attaching. Thus, if one does a simple
+
+       if (openmp_p
+	   && (pragma_kind & GOMP_MAP_VARS_ENTER_DATA)
+	   && mapnum == 1)
+     check in target.c's gomp_map_vars_internal will fail
+     as mapnum > 1 but still the map associated with this
+     ATTACH is in a different set.  */
+
+  struct S s1, *sptr1;
+  s1.ptrset = (int **) __builtin_malloc (sizeof(void*) * 3);
+  s1.ptrset2 = (int **) __builtin_malloc (sizeof(void*) * 3);
+  s1.ptrset[0] = data;
+  s1.ptrset[1] = data;
+  s1.ptrset[2] = data;
+  s1.ptrset2[0] = data;
+  s1.ptrset2[1] = data;
+  s1.ptrset2[2] = data;
+  sptr1 = &s1;
+
+  i = 1;
+  n = 0;
+  n2 = 2;
+  #pragma omp target enter data map(sptr1[:1], sptr1->ptrset[:3], sptr1->ptrset2[:3])
+  #pragma omp target enter data map(sptr1->ptrset[i][:n], sptr1->ptrset2[i][:n])
+
+  #pragma omp target
+    if (sptr1->ptrset2[1][0] != 1 || sptr1->ptrset2[1][1] != 2)
+      __builtin_abort ();
+
+  #pragma omp target exit data map(sptr1->ptrset[i][:n], sptr1->ptrset2[i][:n])
+  #pragma omp target exit data map(sptr1[:1], sptr1->ptrset[:3], sptr1->ptrset2[:3])
+
+  __builtin_free (s1.ptrset);
+  __builtin_free (s1.ptrset2);
+}
diff --git a/libgomp/testsuite/libgomp.c/target-map-zero-sized.c b/libgomp/testsuite/libgomp.c/target-map-zero-sized.c
new file mode 100644
index 0000000..7c4ab80
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-map-zero-sized.c
@@ -0,0 +1,107 @@
+/* { dg-do run } */
+/* { dg-additional-options "-O0" }  */
+
+/* Issue showed up in the real world when large data was distributed
+   over multiple MPI progresses - such that for one process n == 0
+   happend at run time.
+
+   Before  map(var[:0])  and  map(var[:n])  with n > 0 was handled,
+   this patch now also handles  map(var[:n]) with n == 0.
+
+   Failed before with "libgomp: pointer target not mapped for attach".  */
+
+/* Here, the base address is shifted - which should have no effect,
+   but must work as well.  */
+void
+with_offset ()
+{
+  struct S {
+     int *ptr1, *ptr2;
+  };
+  struct S s1, s2;
+  int *a, *b, *c, *d;
+  s1.ptr1 = (int *) 0L;
+  s1.ptr2 = (int *) 0xdeedbeef;
+  s2.ptr1 = (int *) 0L;
+  s2.ptr2 = (int *) 0xdeedbeef;
+  a = (int *) 0L;
+  b = (int *) 0xdeedbeef;
+  c = (int *) 0L;
+  d = (int *) 0xdeedbeef;
+
+  int n1, n2, n3, n4;
+  n1 = n2 = n3 = n4 = 0;
+
+  #pragma omp target enter data map(s1.ptr1[4:n1], s1.ptr2[6:n2], a[3:n3], b[2:n4])
+
+  #pragma omp target map(s2.ptr1[4:n1], s2.ptr2[2:n2], c[6:n3], d[9:n4])
+  {
+    if (s2.ptr1 != (void *) 0L || s2.ptr2 != (void *) 0xdeedbeef
+	|| c != (void *) 0L || d != (void *) 0xdeedbeef)
+      __builtin_abort ();
+  }
+
+  #pragma omp target map(s1.ptr1[4:n1], s1.ptr2[6:n2], a[3:n3], b[2:n4])
+  {
+    if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
+	|| a != (void *) 0L || b != (void *) 0xdeedbeef)
+      __builtin_abort ();
+  }
+
+  #pragma omp target
+  {
+    if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
+	|| a != (void *) 0L || b != (void *) 0xdeedbeef)
+      __builtin_abort ();
+  }
+
+  #pragma omp target exit data map(s1.ptr1[4:n1], s1.ptr2[6:n2], a[3:n3], b[2:n4])
+}
+
+int
+main ()
+{
+  struct S {
+     int *ptr1, *ptr2;
+  };
+  struct S s1, s2;
+  int *a, *b, *c, *d;
+  s1.ptr1 = (int *) 0L;
+  s1.ptr2 = (int *) 0xdeedbeef;
+  s2.ptr1 = (int *) 0L;
+  s2.ptr2 = (int *) 0xdeedbeef;
+  a = (int *) 0L;
+  b = (int *) 0xdeedbeef;
+  c = (int *) 0L;
+  d = (int *) 0xdeedbeef;
+
+  int n1, n2, n3, n4;
+  n1 = n2 = n3 = n4 = 0;
+
+  #pragma omp target enter data map(s1.ptr1[:n1], s1.ptr2[:n2], a[:n3], b[:n4])
+
+  #pragma omp target map(s2.ptr1[:n1], s2.ptr2[:n2], c[:n3], d[:n4])
+  {
+    if (s2.ptr1 != (void *) 0L || s2.ptr2 != (void *) 0xdeedbeef
+	|| c != (void *) 0L || d != (void *) 0xdeedbeef)
+      __builtin_abort ();
+  }
+
+  #pragma omp target map(s1.ptr1[:n1], s1.ptr2[:n2], a[:n3], b[:n4])
+  {
+    if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
+	|| a != (void *) 0L || b != (void *) 0xdeedbeef)
+      __builtin_abort ();
+  }
+
+  #pragma omp target
+  {
+    if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef
+	|| a != (void *) 0L || b != (void *) 0xdeedbeef)
+      __builtin_abort ();
+  }
+
+  #pragma omp target exit data map(s1.ptr1[:n1], s1.ptr2[:n2], a[:n3], b[:n4])
+
+  with_offset ();
+}