OpenMP/OpenACC struct sibling list gimplification extension and rework

This patch refactors struct sibling-list processing in gimplify.cc, and
adjusts some related mapping-clause processing in the Fortran FE and
omp-low.cc accordingly.

2022-09-13  Julian Brown  <julian@codesourcery.com>

gcc/fortran/
	* trans-openmp.cc (gfc_trans_omp_clauses): Don't create
	GOMP_MAP_TO_PSET mappings for class metadata, nor GOMP_MAP_POINTER
	mappings for POINTER_TYPE_P decls.

gcc/
	* gimplify.cc (gimplify_omp_var_data): Remove GOVD_MAP_HAS_ATTACHMENTS.
	(GOMP_FIRSTPRIVATE_IMPLICIT): Renumber.
	(insert_struct_comp_map): Refactor function into...
	(build_omp_struct_comp_nodes): This new function.  Remove list handling
	and improve self-documentation.
	(extract_base_bit_offset): Remove BASE_REF, OFFSETP parameters.  Move
	code to strip outer parts of address out of function, but strip no-op
	conversions.
	(omp_mapping_group): Add DELETED field for use during reindexing.
	(omp_strip_components_and_deref, omp_strip_indirections): New functions.
	(omp_group_last, omp_group_base): Add GOMP_MAP_STRUCT handling.
	(omp_gather_mapping_groups): Initialise DELETED field for new groups.
	(omp_index_mapping_groups): Notice DELETED groups when (re)indexing.
	(omp_siblist_insert_node_after, omp_siblist_move_node_after,
	omp_siblist_move_nodes_after, omp_siblist_move_concat_nodes_after): New
	helper functions.
	(omp_accumulate_sibling_list): New function to build up GOMP_MAP_STRUCT
	node groups for sibling lists. Outlined from gimplify_scan_omp_clauses.
	(omp_build_struct_sibling_lists): New function.
	(gimplify_scan_omp_clauses): Remove struct_map_to_clause,
	struct_seen_clause, struct_deref_set.  Call
	omp_build_struct_sibling_lists as pre-pass instead of handling sibling
	lists in the function's main processing loop.
	(gimplify_adjust_omp_clauses_1): Remove GOVD_MAP_HAS_ATTACHMENTS
	handling, unused now.
	* omp-low.cc (scan_sharing_clauses): Handle pointer-type indirect
	struct references, and references to pointers to structs also.

gcc/testsuite/
	* g++.dg/goacc/member-array-acc.C: New test.
	* g++.dg/gomp/member-array-omp.C: New test.
	* g++.dg/gomp/target-3.C: Update expected output.
	* g++.dg/gomp/target-lambda-1.C: Likewise.
	* g++.dg/gomp/target-this-2.C: Likewise.
	* c-c++-common/goacc/deep-copy-arrayofstruct.c: Move test from here.
	* c-c++-common/gomp/target-50.c: New test.

libgomp/
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c: New test.
	* testsuite/libgomp.oacc-c++/deep-copy-17.C: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c: Move
	test to here, make "run" test.

4 files changed, 483 insertions, 0 deletions

diff --git a/libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C
new file mode 100644
index 0000000..dacbb52
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C
@@ -0,0 +1,101 @@
+#include <cassert>
+
+/* Test attach/detach operation with pointers and references to structs.  */
+
+typedef struct mystruct {
+  int *a;
+  int b;
+  int *c;
+  int d;
+  int *e;
+} mystruct;
+
+void str (void)
+{
+  int a[10], c[10], e[10];
+  mystruct m = { .a = a, .c = c, .e = e };
+  a[0] = 5;
+  c[0] = 7;
+  e[0] = 9;
+  #pragma acc parallel copy(m.a[0:10], m.b, m.c[0:10], m.d, m.e[0:10])
+  {
+    m.a[0] = m.c[0] + m.e[0];
+  }
+  assert (m.a[0] == 7 + 9);
+}
+
+void strp (void)
+{
+  int *a = new int[10];
+  int *c = new int[10];
+  int *e = new int[10];
+  mystruct *m = new mystruct;
+  m->a = a;
+  m->c = c;
+  m->e = e;
+  a[0] = 6;
+  c[0] = 8;
+  e[0] = 10;
+  #pragma acc parallel copy(m->a[0:10], m->b, m->c[0:10], m->d, m->e[0:10])
+  {
+    m->a[0] = m->c[0] + m->e[0];
+  }
+  assert (m->a[0] == 8 + 10);
+  delete m;
+  delete[] a;
+  delete[] c;
+  delete[] e;
+}
+
+void strr (void)
+{
+  int *a = new int[10];
+  int *c = new int[10];
+  int *e = new int[10];
+  mystruct m;
+  mystruct &n = m;
+  n.a = a;
+  n.c = c;
+  n.e = e;
+  a[0] = 7;
+  c[0] = 9;
+  e[0] = 11;
+  #pragma acc parallel copy(n.a[0:10], n.b, n.c[0:10], n.d, n.e[0:10])
+  {
+    n.a[0] = n.c[0] + n.e[0];
+  }
+  assert (n.a[0] == 9 + 11);
+  delete[] a;
+  delete[] c;
+  delete[] e;
+}
+
+void strrp (void)
+{
+  int a[10], c[10], e[10];
+  mystruct *m = new mystruct;
+  mystruct *&n = m;
+  n->a = a;
+  n->b = 3;
+  n->c = c;
+  n->d = 5;
+  n->e = e;
+  a[0] = 8;
+  c[0] = 10;
+  e[0] = 12;
+  #pragma acc parallel copy(n->a[0:10], n->c[0:10], n->e[0:10])
+  {
+    n->a[0] = n->c[0] + n->e[0];
+  }
+  assert (n->a[0] == 10 + 12);
+  delete m;
+}
+
+int main (int argc, char *argv[])
+{
+  str ();
+  strp ();
+  strr ();
+  strrp ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c
new file mode 100644
index 0000000..27fe1a9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c
@@ -0,0 +1,68 @@
+#include <stdlib.h>
+
+/* Test multiple struct dereferences on one directive, and slices starting at
+   non-zero.  */
+
+typedef struct {
+  int *a;
+  int *b;
+  int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct *m = (mystruct *) malloc (sizeof (*m));
+  int i;
+
+  m->a = (int *) malloc (N * sizeof (int));
+  m->b = (int *) malloc (N * sizeof (int));
+  m->c = (int *) malloc (N * sizeof (int));
+
+  for (i = 0; i < N; i++)
+    {
+      m->a[i] = 0;
+      m->b[i] = 0;
+      m->c[i] = 0;
+    }
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+#pragma acc parallel loop copy(m->a[0:N])
+      for (j = 0; j < N; j++)
+	m->a[j]++;
+#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10])
+      for (j = 0; j < N; j++)
+	{
+	  m->b[j]++;
+	  if (j > 5 && j < N - 5)
+	    m->c[j]++;
+	}
+    }
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 99)
+	abort ();
+      if (m->b[i] != 99)
+	abort ();
+      if (i > 5 && i < N-5)
+	{
+	  if (m->c[i] != 99)
+	    abort ();
+	}
+      else
+	{
+	  if (m->c[i] != 0)
+	    abort ();
+	}
+    }
+
+  free (m->a);
+  free (m->b);
+  free (m->c);
+  free (m);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c
new file mode 100644
index 0000000..a7308e8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c
@@ -0,0 +1,231 @@
+#include <stdlib.h>
+
+/* Test mapping chained indirect struct accesses, mixed in different ways.  */
+
+typedef struct {
+  int *a;
+  int b;
+  int *c;
+} str1;
+
+typedef struct {
+  int d;
+  int *e;
+  str1 *f;
+} str2;
+
+typedef struct {
+  int g;
+  int h;
+  str2 *s2;
+} str3;
+
+typedef struct {
+  str3 m;
+  str3 n;
+} str4;
+
+void
+zero_arrays (str4 *s, int N)
+{
+  for (int i = 0; i < N; i++)
+    {
+      s->m.s2->e[i] = 0;
+      s->m.s2->f->a[i] = 0;
+      s->m.s2->f->c[i] = 0;
+      s->n.s2->e[i] = 0;
+      s->n.s2->f->a[i] = 0;
+      s->n.s2->f->c[i] = 0;
+    }
+}
+
+void
+alloc_s2 (str2 **s, int N)
+{
+  (*s) = (str2 *) malloc (sizeof (str2));
+  (*s)->f = (str1 *) malloc (sizeof (str1));
+  (*s)->e = (int *) malloc (sizeof (int) * N);
+  (*s)->f->a = (int *) malloc (sizeof (int) * N);
+  (*s)->f->c = (int *) malloc (sizeof (int) * N);
+}
+
+int main (int argc, char* argv[])
+{
+  const int N = 1024;
+  str4 p, *q;
+  int i;
+
+  alloc_s2 (&p.m.s2, N);
+  alloc_s2 (&p.n.s2, N);
+  q = (str4 *) malloc (sizeof (str4));
+  alloc_s2 (&q->m.s2, N);
+  alloc_s2 (&q->n.s2, N);
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1])
+#pragma acc parallel loop copy(p.m.s2->e[:N])
+      for (int j = 0; j < N; j++)
+	p.m.s2->e[j]++;
+#pragma acc exit data delete(p.m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->e[i] != 99)
+      abort ();
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1])
+#pragma acc enter data copyin(p.m.s2->f[:1])
+#pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.m.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    p.m.s2->f->a[j]++;
+	    p.m.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(p.m.s2->f[:1])
+#pragma acc exit data delete(p.m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->f->a[i] != 99 || p.m.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1]) copyin(p.n.s2[:1])
+#pragma acc enter data copyin(p.m.s2->f[:1]) copyin(p.n.s2->f[:1])
+#pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.m.s2->f->c[:N]) \
+			  copy(p.n.s2->f->a[:N]) copy(p.n.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    p.m.s2->f->a[j]++;
+	    p.m.s2->f->c[j]++;
+	    p.n.s2->f->a[j]++;
+	    p.n.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(p.m.s2->f[:1]) delete(p.n.s2->f[:1])
+#pragma acc exit data delete(p.m.s2[:1]) delete(p.n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->f->a[i] != 99 || p.m.s2->f->c[i] != 99
+	|| p.n.s2->f->a[i] != 99 || p.n.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1]) copyin(p.n.s2[:1])
+#pragma acc enter data copyin(p.n.s2->e[:N]) copyin(p.n.s2->f[:1]) \
+		       copyin(p.m.s2->f[:1])
+#pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.n.s2->f->a[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    p.m.s2->f->a[j]++;
+	    p.n.s2->f->a[j]++;
+	    p.n.s2->e[j]++;
+	  }
+#pragma acc exit data delete(p.m.s2->f[:1]) delete(p.n.s2->f[:1]) \
+		      copyout(p.n.s2->e[:N])
+#pragma acc exit data delete(p.m.s2[:1]) delete(p.n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->f->a[i] != 99 || p.n.s2->f->a[i] != 99
+	|| p.n.s2->e[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1])
+#pragma acc parallel loop copy(q->m.s2->e[:N])
+      for (int j = 0; j < N; j++)
+	q->m.s2->e[j]++;
+#pragma acc exit data delete(q->m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->e[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1])
+#pragma acc enter data copyin(q->m.s2->f[:1])
+#pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->m.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    q->m.s2->f->a[j]++;
+	    q->m.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(q->m.s2->f[:1])
+#pragma acc exit data delete(q->m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->f->a[i] != 99 || q->m.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1]) copyin(q->n.s2[:1])
+#pragma acc enter data copyin(q->m.s2->f[:1]) copyin(q->n.s2->f[:1])
+#pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->m.s2->f->c[:N]) \
+			  copy(q->n.s2->f->a[:N]) copy(q->n.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    q->m.s2->f->a[j]++;
+	    q->m.s2->f->c[j]++;
+	    q->n.s2->f->a[j]++;
+	    q->n.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(q->m.s2->f[:1]) delete(q->n.s2->f[:1])
+#pragma acc exit data delete(q->m.s2[:1]) delete(q->n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->f->a[i] != 99 || q->m.s2->f->c[i] != 99
+	|| q->n.s2->f->a[i] != 99 || q->n.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1]) copyin(q->n.s2[:1])
+#pragma acc enter data copyin(q->n.s2->e[:N]) copyin(q->m.s2->f[:1]) \
+		       copyin(q->n.s2->f[:1])
+#pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->n.s2->f->a[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    q->m.s2->f->a[j]++;
+	    q->n.s2->f->a[j]++;
+	    q->n.s2->e[j]++;
+	  }
+#pragma acc exit data delete(q->m.s2->f[:1]) delete(q->n.s2->f[:1]) \
+		      copyout(q->n.s2->e[:N])
+#pragma acc exit data delete(q->m.s2[:1]) delete(q->n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->f->a[i] != 99 || q->n.s2->f->a[i] != 99
+	|| q->n.s2->e[i] != 99)
+      abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c
new file mode 100644
index 0000000..a11c647
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c
@@ -0,0 +1,83 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <stdio.h>
+
+typedef struct {
+  int *a;
+  int *b;
+  int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+  const int N = 1024;
+  const int S = 32;
+  mystruct *m = (mystruct *) calloc (S, sizeof (*m));
+  int i, j;
+
+  for (i = 0; i < S; i++)
+    {
+      m[i].a = (int *) malloc (N * sizeof (int));
+      m[i].b = (int *) malloc (N * sizeof (int));
+      m[i].c = (int *) malloc (N * sizeof (int));
+    }
+
+  for (j = 0; j < S; j++)
+    for (i = 0; i < N; i++)
+      {
+	m[j].a[i] = 0;
+	m[j].b[i] = 0;
+	m[j].c[i] = 0;
+      }
+
+#pragma acc enter data copyin(m[0:1])
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j, k;
+      for (k = 0; k < S; k++)
+#pragma acc parallel loop copy(m[k].a[0:N])
+        for (j = 0; j < N; j++)
+          m[k].a[j]++;
+
+      for (k = 0; k < S; k++)
+#pragma acc parallel loop copy(m[k].b[0:N], m[k].c[5:N-10])
+	for (j = 0; j < N; j++)
+	  {
+	    m[k].b[j]++;
+	    if (j > 5 && j < N - 5)
+	      m[k].c[j]++;
+	}
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  for (j = 0; j < S; j++)
+    {
+      for (i = 0; i < N; i++)
+	{
+	  if (m[j].a[i] != 99)
+	    abort ();
+	  if (m[j].b[i] != 99)
+	    abort ();
+	  if (i > 5 && i < N-5)
+	    {
+	      if (m[j].c[i] != 99)
+		abort ();
+	    }
+	  else
+	    {
+	      if (m[j].c[i] != 0)
+		abort ();
+	    }
+	}
+
+      free (m[j].a);
+      free (m[j].b);
+      free (m[j].c);
+    }
+  free (m);
+
+  return 0;
+}