OpenMP/OpenACC: Unordered/non-constant component offset runtime diagnostic

This patch adds support for non-constant component offsets in "map"
clauses for OpenMP (and the equivalants for OpenACC), which are not able
to be sorted into order at compile time.  Normally struct accesses in
such clauses are gathered together and sorted into increasing address
order after a "GOMP_MAP_STRUCT" node: if we have variable indices,
that is no longer possible.

This version of the patch scales back the previously-posted version to
merely add a diagnostic for incorrect usage of component accesses with
variably-indexed arrays of structs: the only permitted variant is where
we have multiple indices that are the same, but we could not prove so
at compile time.  Rather than silently producing the wrong result for
cases where the indices are in fact different, we error out (e.g.,
"map(dtarr(i)%arrptr, dtarr(j)%arrptr(4:8))", for different i/j).

For now, multiple *constant* array indices are still supported (see
map-arrayofstruct-1.c).  That could perhaps be addressed with a follow-up
patch, if necessary.

This version of the patch renumbers the GOMP_MAP_STRUCT_UNORD kind to
avoid clashing with the OpenACC "non-contiguous" dynamic array support
(though that is not yet applied to mainline).

2023-08-18  Julian Brown  <julian@codesourcery.com>

gcc/
	* gimplify.cc (extract_base_bit_offset): Add VARIABLE_OFFSET parameter.
	(omp_get_attachment, omp_group_last, omp_group_base,
	omp_directive_maps_explicitly): Add GOMP_MAP_STRUCT_UNORD support.
	(omp_accumulate_sibling_list): Update calls to extract_base_bit_offset.
	Support GOMP_MAP_STRUCT_UNORD.
	(omp_build_struct_sibling_lists, gimplify_scan_omp_clauses,
	gimplify_adjust_omp_clauses, gimplify_omp_target_update): Add
	GOMP_MAP_STRUCT_UNORD support.
	* omp-low.cc (lower_omp_target): Add GOMP_MAP_STRUCT_UNORD support.
	* tree-pretty-print.cc (dump_omp_clause): Likewise.

include/
	* gomp-constants.h (gomp_map_kind): Add GOMP_MAP_STRUCT_UNORD.

libgomp/
	* oacc-mem.c (find_group_last, goacc_enter_data_internal,
	goacc_exit_data_internal, GOACC_enter_exit_data): Add
	GOMP_MAP_STRUCT_UNORD support.
	* target.c (gomp_map_vars_internal): Add GOMP_MAP_STRUCT_UNORD support.
	Detect incorrect use of variable indexing of arrays of structs.
	(GOMP_target_enter_exit_data, gomp_target_task_fn): Add
	GOMP_MAP_STRUCT_UNORD support.
	* testsuite/libgomp.c-c++-common/map-arrayofstruct-1.c: New test.
	* testsuite/libgomp.c-c++-common/map-arrayofstruct-2.c: New test.
	* testsuite/libgomp.c-c++-common/map-arrayofstruct-3.c: New test.
	* testsuite/libgomp.fortran/map-subarray-5.f90: New test.

6 files changed, 290 insertions, 7 deletions

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index fe63274..79d6e32 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1028,6 +1028,7 @@ find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds)
       break;
 
     case GOMP_MAP_STRUCT:
+    case GOMP_MAP_STRUCT_UNORD:
       pos += sizes[pos];
       break;
 
@@ -1088,6 +1089,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
       switch (kinds[i] & 0xff)
 	{
 	case GOMP_MAP_STRUCT:
+	case GOMP_MAP_STRUCT_UNORD:
 	  {
 	    size = (uintptr_t) hostaddrs[group_last] + sizes[group_last]
 		   - (uintptr_t) hostaddrs[i];
@@ -1334,6 +1336,7 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 	  break;
 
 	case GOMP_MAP_STRUCT:
+	case GOMP_MAP_STRUCT_UNORD:
 	  /* Skip the 'GOMP_MAP_STRUCT' itself, and use the regular processing
 	     for all its entries.  This special handling exists for GCC 10.1
 	     compatibility; afterwards, we're not generating these no-op
@@ -1472,7 +1475,8 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs,
 
       if (kind == GOMP_MAP_POINTER
 	  || kind == GOMP_MAP_TO_PSET
-	  || kind == GOMP_MAP_STRUCT)
+	  || kind == GOMP_MAP_STRUCT
+	  || kind == GOMP_MAP_STRUCT_UNORD)
 	continue;
 
       if (kind == GOMP_MAP_FORCE_ALLOC
diff --git a/libgomp/target.c b/libgomp/target.c
index f435521..434d40d 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1084,7 +1084,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 	    tgt->list[i].offset = 0;
 	  continue;
 	}
-      else if ((kind & typemask) == GOMP_MAP_STRUCT)
+      else if ((kind & typemask) == GOMP_MAP_STRUCT
+	       || (kind & typemask) == GOMP_MAP_STRUCT_UNORD)
 	{
 	  size_t first = i + 1;
 	  size_t last = i + sizes[i];
@@ -1468,6 +1469,20 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 		    tgt->list[i].offset = OFFSET_INLINED;
 		  }
 		continue;
+	      case GOMP_MAP_STRUCT_UNORD:
+		if (sizes[i] > 1)
+		  {
+		    void *first = hostaddrs[i + 1];
+		    for (size_t j = i + 1; j < i + sizes[i]; j++)
+		      if (hostaddrs[j + 1] != first)
+			{
+			  gomp_mutex_unlock (&devicep->lock);
+			  gomp_fatal ("Mapped array elements must be the "
+				      "same (%p vs %p)", first,
+				      hostaddrs[j + 1]);
+			}
+		  }
+		/* Fallthrough.  */
 	      case GOMP_MAP_STRUCT:
 		first = i + 1;
 		last = i + sizes[i];
@@ -1586,9 +1601,40 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 	      k->host_end = k->host_start + sizeof (void *);
 	    splay_tree_key n = splay_tree_lookup (mem_map, k);
 	    if (n && n->refcount != REFCOUNT_LINK)
-	      gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i],
-				      kind & typemask, false, implicit, cbufp,
-				      refcount_set);
+	      {
+		if (field_tgt_clear != FIELD_TGT_EMPTY)
+		  {
+		    /* For this condition to be true, there must be a
+		       duplicate struct element mapping.  This can happen with
+		       GOMP_MAP_STRUCT_UNORD mappings, for example.  */
+		    tgt->list[i].key = n;
+		    if (openmp_p)
+		      {
+			assert ((n->refcount & REFCOUNT_STRUCTELEM) != 0);
+			assert (field_tgt_structelem_first != NULL);
+
+			if (i == field_tgt_clear)
+			  {
+			    n->refcount |= REFCOUNT_STRUCTELEM_FLAG_LAST;
+			    field_tgt_structelem_first = NULL;
+			  }
+		      }
+		    if (i == field_tgt_clear)
+		      field_tgt_clear = FIELD_TGT_EMPTY;
+		    gomp_increment_refcount (n, refcount_set);
+		    tgt->list[i].copy_from
+		      = GOMP_MAP_COPY_FROM_P (kind & typemask);
+		    tgt->list[i].always_copy_from
+		      = GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
+		    tgt->list[i].is_attach = false;
+		    tgt->list[i].offset = 0;
+		    tgt->list[i].length = k->host_end - k->host_start;
+		  }
+		else
+		  gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i],
+					  kind & typemask, false, implicit,
+					  cbufp, refcount_set);
+	      }
 	    else
 	      {
 		k->aux = NULL;
@@ -4166,7 +4212,8 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
   size_t i, j;
   if ((flags & GOMP_TARGET_FLAG_EXIT_DATA) == 0)
     for (i = 0; i < mapnum; i++)
-      if ((kinds[i] & 0xff) == GOMP_MAP_STRUCT)
+      if ((kinds[i] & 0xff) == GOMP_MAP_STRUCT
+	  || (kinds[i] & 0xff) == GOMP_MAP_STRUCT_UNORD)
 	{
 	  gomp_map_vars (devicep, sizes[i] + 1, &hostaddrs[i], NULL, &sizes[i],
 			 &kinds[i], true, &refcount_set,
@@ -4264,7 +4311,8 @@ gomp_target_task_fn (void *data)
       htab_t refcount_set = htab_create (ttask->mapnum);
       if ((ttask->flags & GOMP_TARGET_FLAG_EXIT_DATA) == 0)
 	for (i = 0; i < ttask->mapnum; i++)
-	  if ((ttask->kinds[i] & 0xff) == GOMP_MAP_STRUCT)
+	  if ((ttask->kinds[i] & 0xff) == GOMP_MAP_STRUCT
+	      || (ttask->kinds[i] & 0xff) == GOMP_MAP_STRUCT_UNORD)
 	    {
 	      gomp_map_vars (devicep, ttask->sizes[i] + 1, &ttask->hostaddrs[i],
 			     NULL, &ttask->sizes[i], &ttask->kinds[i], true,
diff --git a/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-1.c b/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-1.c
new file mode 100644
index 0000000..655f6ef
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-1.c
@@ -0,0 +1,51 @@
+#include <stdlib.h>
+#include <assert.h>
+
+struct st {
+  int *p;
+};
+
+int main (void)
+{
+  struct st s[2];
+  s[0].p = (int *) calloc (5, sizeof (int));
+  s[1].p = (int *) calloc (5, sizeof (int));
+
+  /* These mappings not supported by the OpenMP spec, and are currently
+     implemented as an extension by GCC for legacy compatibility only.  See
+     e.g. OpenMP 5.2, "5.8.3 map Clause":
+
+    "If multiple list items are explicitly mapped on the same construct and
+     have the same containing array or have base pointers that share original
+     storage, and if any of the list items do not have corresponding list
+     items that are present in the device data environment prior to a task
+     encountering the construct, then the list items must refer to the same
+     array elements of either the containing array or the implicit array of
+     the base pointers."
+  */
+
+#pragma omp target map(s[0].p, s[1].p, s[0].p[0:2], s[1].p[1:3])
+  {
+    s[0].p[0] = 5;
+    s[1].p[1] = 7;
+  }
+
+#pragma omp target map(s, s[0].p[0:2], s[1].p[1:3])
+  {
+    s[0].p[0]++;
+    s[1].p[1]++;
+  }
+
+#pragma omp target map(s[0:2], s[0].p[0:2], s[1].p[1:3])
+  {
+    s[0].p[0]++;
+    s[1].p[1]++;
+  }
+
+  assert (s[0].p[0] == 7);
+  assert (s[1].p[1] == 9);
+
+  free (s[0].p);
+  free (s[1].p);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-2.c b/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-2.c
new file mode 100644
index 0000000..ff7ce0e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-2.c
@@ -0,0 +1,58 @@
+#include <stdlib.h>
+#include <assert.h>
+
+struct st {
+  int *p;
+};
+
+int main (void)
+{
+  struct st s[10];
+
+  for (int i = 0; i < 10; i++)
+    s[i].p = (int *) calloc (5, sizeof (int));
+
+  for (int i = 0; i < 10; i++)
+    for (int j = 0; j < 10; j++)
+      for (int k = 0; k < 10; k++)
+	{
+	  if (i == j || j == k || i == k)
+	    continue;
+
+#pragma omp target map(s[i].p, s[j].p, s[k].p, s[i].p[0:2], s[j].p[1:3], \
+		       s[k].p[2])
+	  {
+	    s[i].p[0]++;
+	    s[j].p[1]++;
+	    s[k].p[2]++;
+	  }
+
+#pragma omp target map(s, s[i].p[0:2], s[j].p[1:3], s[k].p[2])
+	  {
+	    s[i].p[0]++;
+	    s[j].p[1]++;
+	    s[k].p[2]++;
+	  }
+
+#pragma omp target map(s[0:10], s[i].p[0:2], s[j].p[1:3], s[k].p[2])
+	  {
+	    s[i].p[0]++;
+	    s[j].p[1]++;
+	    s[k].p[2]++;
+	  }
+	}
+
+  for (int i = 0; i < 10; i++)
+    {
+      assert (s[i].p[0] == 216);
+      assert (s[i].p[1] == 216);
+      assert (s[i].p[2] == 216);
+      free (s[i].p);
+    }
+
+  return 0;
+}
+
+/* { dg-output "(\n|\r|\r\n)" { target offload_device_nonshared_as } } */
+/* { dg-output "libgomp: Mapped array elements must be the same .*(\n|\r|\r\n)+" { target offload_device_nonshared_as } } */
+/* { dg-shouldfail "" { offload_device_nonshared_as } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-3.c b/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-3.c
new file mode 100644
index 0000000..770ac2a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/map-arrayofstruct-3.c
@@ -0,0 +1,68 @@
+#include <stdlib.h>
+#include <assert.h>
+
+struct st {
+  int *p;
+};
+
+struct tt {
+  struct st a[10];
+};
+
+struct ut {
+  struct tt *t;
+};
+
+int main (void)
+{
+  struct tt *t = (struct tt *) malloc (sizeof *t);
+  struct ut *u = (struct ut *) malloc (sizeof *u);
+
+  for (int i = 0; i < 10; i++)
+    t->a[i].p = (int *) calloc (5, sizeof (int));
+
+  u->t = t;
+
+  for (int i = 0; i < 10; i++)
+    for (int j = 0; j < 10; j++)
+      for (int k = 0; k < 10; k++)
+	{
+	  if (i == j || j == k || i == k)
+	    continue;
+
+	  /* This one can use "firstprivate" for T...  */
+#pragma omp target map(t->a[i].p, t->a[j].p, t->a[k].p, \
+		       t->a[i].p[0:2], t->a[j].p[1:3], t->a[k].p[2])
+	  {
+	    t->a[i].p[0]++;
+	    t->a[j].p[1]++;
+	    t->a[k].p[2]++;
+	  }
+
+	  /* ...but this one must use attach/detach for T.  */
+#pragma omp target map(u->t, u->t->a[i].p, u->t->a[j].p, u->t->a[k].p, \
+		       u->t->a[i].p[0:2], u->t->a[j].p[1:3], u->t->a[k].p[2])
+	  {
+	    u->t->a[i].p[0]++;
+	    u->t->a[j].p[1]++;
+	    u->t->a[k].p[2]++;
+	  }
+	}
+
+  for (int i = 0; i < 10; i++)
+    {
+      assert (t->a[i].p[0] == 144);
+      assert (t->a[i].p[1] == 144);
+      assert (t->a[i].p[2] == 144);
+      free (t->a[i].p);
+    }
+
+  free (u);
+  free (t);
+
+  return 0;
+}
+
+/* { dg-output "(\n|\r|\r\n)" { target offload_device_nonshared_as } } */
+/* { dg-output "libgomp: Mapped array elements must be the same .*(\n|\r|\r\n)+" { target offload_device_nonshared_as } } */
+/* { dg-shouldfail "" { offload_device_nonshared_as } } */
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-5.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-5.f90
new file mode 100644
index 0000000..e7cdf11
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-5.f90
@@ -0,0 +1,54 @@
+! { dg-do run }
+
+type t
+  integer, pointer :: p(:)
+end type t
+
+type(t) :: var(3)
+integer :: i, j
+
+allocate (var(1)%p, source=[1,2,3,5])
+allocate (var(2)%p, source=[2,3,5])
+allocate (var(3)%p(1:3))
+
+var(3)%p = 0
+
+do i = 1, 3
+  do j = 1, 3
+!$omp target map(var(i)%p, var(j)%p)
+    var(i)%p(1) = 5
+    var(j)%p(2) = 7
+!$omp end target
+
+    if (i.ne.j) then
+!$omp target map(var(i)%p(1:3), var(i)%p, var(j)%p)
+      var(i)%p(1) = var(i)%p(1) + 1
+      var(j)%p(2) = var(j)%p(2) + 1
+!$omp end target
+
+!$omp target map(var(i)%p, var(j)%p, var(j)%p(1:3))
+      var(i)%p(1) = var(i)%p(1) + 1
+      var(j)%p(2) = var(j)%p(2) + 1
+!$omp end target
+
+!$omp target map(var(i)%p, var(i)%p(1:3), var(j)%p, var(j)%p(2))
+      var(i)%p(1) = var(i)%p(1) + 1
+      var(j)%p(2) = var(j)%p(2) + 1
+!$omp end target
+    end if
+
+    if (i.eq.j) then
+      if (var(i)%p(1).ne.5) stop 1
+      if (var(j)%p(2).ne.7) stop 2
+    else
+      if (var(i)%p(1).ne.8) stop 3
+      if (var(j)%p(2).ne.10) stop 4
+    end if
+  end do
+end do
+
+end
+
+! { dg-output "(\n|\r|\r\n)" }
+! { dg-output "libgomp: Mapped array elements must be the same .*(\n|\r|\r\n)+" }
+! { dg-shouldfail "" { offload_device_nonshared_as } }