[PATCH v2 05/11] OpenMP: Handle reference-typed struct members

[Julian Brown <julian@codesourcery.com>]

This patch relates to OpenMP mapping clauses containing struct members of
reference type, e.g. "mystruct.myref.myptr[:N]".  To be able to access
the array slice through the reference in the middle, we need to perform
an attach action for that reference, since it is represented internally
as a pointer.

I don't think the spec allows for this case explicitly.  The closest
clause is (OpenMP 5.0, "2.19.7.1 map Clause"):

  "If the type of a list item is a reference to a type T then the
   reference in the device data environment is initialized to refer to
   the object in the device data environment that corresponds to the
   object referenced by the list item. If mapping occurs, it occurs as
   though the object were mapped through a pointer with an array section
   of type T and length one."

The patch as is allows the mapping to work with just
"mystruct.myref.myptr[:N]", without an explicit "mystruct.myref"
mapping also (because, would that refer to the hidden pointer used by
the reference, or the automatically-dereferenced data itself?). An
attach/detach operation is thus synthesised for the reference.

Reworking the previous "address inspector" patch, it occurred to me that
this patch might be better implemented in the frontend (or frontends).
I've added a note to that effect, but not actually made that change for
the time being.
---
 gcc/gimplify.cc                            |  59 ++++-
 libgomp/testsuite/libgomp.c++/baseptrs-3.C | 275 +++++++++++++++++++++
 2 files changed, 327 insertions(+), 7 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.c++/baseptrs-3.C

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 598c65eb430..0b8c221e4c8 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9813,7 +9813,10 @@ accumulate_sibling_list (enum omp_region_type region_type, enum tree_code code,
   /* FIXME: If we're not mapping the base pointer in some other clause on this
      directive, I think we want to create ALLOC/RELEASE here -- i.e. not
      early-exit.  */
-  if (openmp && attach_detach)
+  if (openmp
+      && attach_detach
+      && !(TREE_CODE (TREE_TYPE (ocd)) == REFERENCE_TYPE
+	   && TREE_CODE (TREE_TYPE (TREE_TYPE (ocd))) != POINTER_TYPE))
     return NULL;
 
   if (!struct_map_to_clause || struct_map_to_clause->get (base) == NULL)
@@ -9862,9 +9865,37 @@ accumulate_sibling_list (enum omp_region_type region_type, enum tree_code code,
 
       tree noind = strip_indirections (base);
 
-      if (!openmp
+      /* TODO: the following two stanzas handling reference-typed struct
+	 members (for OpenMP) and nested base pointers (for OpenACC) could
+	 probably both be better handled in the frontends.  Doing that would
+	 avoid this late manipulation of the clause list.  */
+
+      if (openmp
+	  && TREE_CODE (TREE_TYPE (noind)) == REFERENCE_TYPE
 	  && (region_type & ORT_TARGET)
 	  && TREE_CODE (noind) == COMPONENT_REF)
+	{
+	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end),
+				      OMP_CLAUSE_MAP);
+	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_TO);
+	  OMP_CLAUSE_DECL (c2) = unshare_expr (base);
+	  OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (TREE_TYPE (noind));
+
+	  tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end),
+				      OMP_CLAUSE_MAP);
+	  OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH_DETACH);
+	  OMP_CLAUSE_DECL (c3) = unshare_expr (noind);
+	  OMP_CLAUSE_SIZE (c3) = size_zero_node;
+
+	  OMP_CLAUSE_CHAIN (c2) = c3;
+	  OMP_CLAUSE_CHAIN (c3) = NULL_TREE;
+
+	  *inner = c2;
+	  return NULL;
+	}
+      else if (!openmp
+	       && (region_type & ORT_TARGET)
+	       && TREE_CODE (noind) == COMPONENT_REF)
 	{
 	  /* The base for this component access is a struct component access
 	     itself.  Insert a node to be processed on the next iteration of
@@ -9878,13 +9909,28 @@ accumulate_sibling_list (enum omp_region_type region_type, enum tree_code code,
 	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FORCE_PRESENT);
 	  OMP_CLAUSE_DECL (c2) = unshare_expr (noind);
 	  OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (TREE_TYPE (noind));
+	  OMP_CLAUSE_CHAIN (c2) = NULL_TREE;
 	  *inner = c2;
 	  return NULL;
 	}
 
-      tree sdecl = strip_components_and_deref (base);
+      tree sdecl = base;
+      if (TREE_CODE (sdecl) == INDIRECT_REF
+	  || TREE_CODE (sdecl) == MEM_REF)
+	{
+	  sdecl = TREE_OPERAND (sdecl, 0);
+	  if (TREE_CODE (sdecl) == INDIRECT_REF
+	      && (TREE_CODE (TREE_TYPE (TREE_OPERAND (sdecl, 0)))
+		  == REFERENCE_TYPE))
+	    sdecl = TREE_OPERAND (sdecl, 0);
+	}
 
-      if (POINTER_TYPE_P (TREE_TYPE (sdecl)) && (region_type & ORT_TARGET))
+      while (TREE_CODE (sdecl) == POINTER_PLUS_EXPR)
+	sdecl = TREE_OPERAND (sdecl, 0);
+
+      if (DECL_P (sdecl)
+	  && POINTER_TYPE_P (TREE_TYPE (sdecl))
+	  && (region_type & ORT_TARGET))
 	{
 	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end),
 				      OMP_CLAUSE_MAP);
@@ -10228,11 +10274,10 @@ omp_build_struct_sibling_lists (enum tree_code code,
 	      else
 		*tail = inner;
 
-	      OMP_CLAUSE_CHAIN (inner) = NULL_TREE;
-
 	      omp_mapping_group newgrp;
 	      newgrp.grp_start = new_next ? new_next : tail;
-	      newgrp.grp_end = inner;
+	      newgrp.grp_end = (OMP_CLAUSE_CHAIN (inner)
+				? OMP_CLAUSE_CHAIN (inner) : inner);
 	      newgrp.mark = UNVISITED;
 	      newgrp.sibling = NULL;
 	      newgrp.deleted = false;
diff --git a/libgomp/testsuite/libgomp.c++/baseptrs-3.C b/libgomp/testsuite/libgomp.c++/baseptrs-3.C
new file mode 100644
index 00000000000..39a48a40920
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/baseptrs-3.C
@@ -0,0 +1,275 @@
+#include <cstdlib>
+#include <cstring>
+#include <cassert>
+
+struct sa0
+{
+  int *ptr;
+};
+
+struct sb0
+{
+  int arr[10];
+};
+
+struct sc0
+{
+  sa0 a;
+  sb0 b;
+  sc0 (sa0 &my_a, sb0 &my_b) : a(my_a), b(my_b) {}
+};
+
+void
+foo0 ()
+{
+  sa0 my_a;
+  sb0 my_b;
+
+  my_a.ptr = (int *) malloc (sizeof (int) * 10);
+  sc0 my_c(my_a, my_b);
+
+  memset (my_c.a.ptr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_c.a.ptr, my_c.a.ptr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_c.a.ptr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_c.a.ptr[i] == i);
+
+  memset (my_c.b.arr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_c.b.arr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_c.b.arr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_c.b.arr[i] == i);
+
+  free (my_a.ptr);
+}
+
+struct sa
+{
+  int *ptr;
+};
+
+struct sb
+{
+  int arr[10];
+};
+
+struct sc
+{
+  sa &a;
+  sb &b;
+  sc (sa &my_a, sb &my_b) : a(my_a), b(my_b) {}
+};
+
+void
+foo ()
+{
+  sa my_a;
+  sb my_b;
+
+  my_a.ptr = (int *) malloc (sizeof (int) * 10);
+  sc my_c(my_a, my_b);
+
+  memset (my_c.a.ptr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_c.a.ptr, my_c.a.ptr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_c.a.ptr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_c.a.ptr[i] == i);
+
+  memset (my_c.b.arr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_c.b.arr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_c.b.arr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_c.b.arr[i] == i);
+
+  free (my_a.ptr);
+}
+
+void
+bar ()
+{
+  sa my_a;
+  sb my_b;
+
+  my_a.ptr = (int *) malloc (sizeof (int) * 10);
+  sc my_c(my_a, my_b);
+  sc &my_cref = my_c;
+
+  memset (my_cref.a.ptr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_cref.a.ptr, my_cref.a.ptr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_cref.a.ptr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_cref.a.ptr[i] == i);
+
+  memset (my_cref.b.arr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_cref.b.arr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_cref.b.arr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_cref.b.arr[i] == i);
+
+  free (my_a.ptr);
+}
+
+struct scp0
+{
+  sa *a;
+  sb *b;
+  scp0 (sa *my_a, sb *my_b) : a(my_a), b(my_b) {}
+};
+
+void
+foop0 ()
+{
+  sa *my_a = new sa;
+  sb *my_b = new sb;
+
+  my_a->ptr = new int[10];
+  scp0 *my_c = new scp0(my_a, my_b);
+
+  memset (my_c->a->ptr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_c->a, my_c->a[:1], my_c->a->ptr, my_c->a->ptr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_c->a->ptr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_c->a->ptr[i] == i);
+
+  memset (my_c->b->arr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_c->b, my_c->b[:1], my_c->b->arr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_c->b->arr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_c->b->arr[i] == i);
+
+  delete[] my_a->ptr;
+  delete my_a;
+  delete my_b;
+}
+
+struct scp
+{
+  sa *&a;
+  sb *&b;
+  scp (sa *&my_a, sb *&my_b) : a(my_a), b(my_b) {}
+};
+
+void
+foop ()
+{
+  sa *my_a = new sa;
+  sb *my_b = new sb;
+
+  my_a->ptr = new int[10];
+  scp *my_c = new scp(my_a, my_b);
+
+  memset (my_c->a->ptr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_c->a, my_c->a[:1], my_c->a->ptr, my_c->a->ptr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_c->a->ptr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_c->a->ptr[i] == i);
+
+  memset (my_c->b->arr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_c->b, my_c->b[:1], my_c->b->arr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_c->b->arr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_c->b->arr[i] == i);
+
+  delete[] my_a->ptr;
+  delete my_a;
+  delete my_b;
+}
+
+void
+barp ()
+{
+  sa *my_a = new sa;
+  sb *my_b = new sb;
+
+  my_a->ptr = new int[10];
+  scp *my_c = new scp(my_a, my_b);
+  scp *&my_cref = my_c;
+
+  memset (my_cref->a->ptr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_cref->a, my_cref->a[:1], my_cref->a->ptr, \
+			  my_cref->a->ptr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_cref->a->ptr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_cref->a->ptr[i] == i);
+
+  memset (my_cref->b->arr, 0, sizeof (int) * 10);
+
+  #pragma omp target map (my_cref->b, my_cref->b[:1], my_cref->b->arr[:10])
+  {
+    for (int i = 0; i < 10; i++)
+      my_cref->b->arr[i] = i;
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (my_cref->b->arr[i] == i);
+
+  delete my_a->ptr;
+  delete my_a;
+  delete my_b;
+}
+
+int main (int argc, char *argv[])
+{
+  foo0 ();
+  foo ();
+  bar ();
+  foop0 ();
+  foop ();
+  barp ();
+  return 0;
+}
-- 
2.29.2