[10/16] OpenMP: Fix non-zero attach/detach bias for struct dereferences

Message ID 20211125141013.113782-1-julian@codesourcery.com
State New
Headers show
Series
  • OpenMP: lvalues in "map" clauses and struct handling rework
Related show

Commit Message

Julian Brown Nov. 25, 2021, 2:10 p.m.
This patch fixes attach/detach operations for OpenMP that have a non-zero
bias: these can occur if we have a mapping such as:

  #pragma omp target map(mystruct->a.b[idx].c[:arrsz])

i.e. where there is an offset between the attachment point ("mystruct"
here) and the pointed-to data.  (The "b" and "c" members would be array
types here, not pointers themselves).  In this example the difference
(thus bias encoded in the attach/detach node) will be something like:

  (uintptr_t) &mystruct->a.b[idx].c[0] - (uintptr_t) &mystruct->a

OK?

Thanks,

Julian

2021-09-29  Julian Brown  <julian@codesourcery.com>

gcc/c-family/
	* c-common.h (c_omp_decompose_attachable_address): Add prototype.
	* c-omp.c (c_omp_decompose_attachable_address): New function.

gcc/c/
	* c-typeck.c (handle_omp_array_sections): Handle attach/detach for
	struct dereferences with non-zero bias.

gcc/cp/
	* semantics.c (handle_omp_array_section): Handle attach/detach for
	struct dereferences with non-zero bias.

libgomp/
	* testsuite/libgomp.c++/baseptrs-3.C: Add test (XFAILed for now).
	* testsuite/libgomp.c-c++-common/baseptrs-1.c: Add test.
	* testsuite/libgomp.c-c++-common/baseptrs-2.c: Add test.
---
 gcc/c-family/c-common.h                       |   1 +
 gcc/c-family/c-omp.c                          |  42 ++++
 gcc/c/c-typeck.c                              |  12 +-
 gcc/cp/semantics.c                            |  14 +-
 libgomp/testsuite/libgomp.c++/baseptrs-3.C    | 182 ++++++++++++++++++
 .../libgomp.c-c++-common/baseptrs-1.c         |  50 +++++
 .../libgomp.c-c++-common/baseptrs-2.c         |  70 +++++++
 7 files changed, 364 insertions(+), 7 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.c++/baseptrs-3.C
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/baseptrs-1.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/baseptrs-2.c

-- 
2.29.2

Patch

diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index d5dad99ff97..dd103d8eecd 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1251,6 +1251,7 @@  extern tree c_omp_check_context_selector (location_t, tree);
 extern void c_omp_mark_declare_variant (location_t, tree, tree);
 extern const char *c_omp_map_clause_name (tree, bool);
 extern void c_omp_adjust_map_clauses (tree, bool);
+extern tree c_omp_decompose_attachable_address (tree t, tree *virtbase);
 
 enum c_omp_directive_kind {
   C_OMP_DIR_STANDALONE,
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index 3f84fd1b5cb..a90696fe706 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -3113,6 +3113,48 @@  c_omp_adjust_map_clauses (tree clauses, bool is_target)
     }
 }
 
+tree
+c_omp_decompose_attachable_address (tree t, tree *virtbase)
+{
+  *virtbase = t;
+
+  /* It's already a pointer.  Just use that.  */
+  if (POINTER_TYPE_P (TREE_TYPE (t)))
+    return NULL_TREE;
+
+  /* Otherwise, look for a base pointer deeper within the expression.  */
+
+  while (TREE_CODE (t) == COMPONENT_REF
+	 && (TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF
+	     || TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF))
+    {
+      t = TREE_OPERAND (t, 0);
+      while (TREE_CODE (t) == ARRAY_REF)
+	t = TREE_OPERAND (t, 0);
+    }
+
+
+  *virtbase = t;
+
+  if (TREE_CODE (t) != COMPONENT_REF)
+    return NULL_TREE;
+
+  t = TREE_OPERAND (t, 0);
+
+  tree attach_pt = NULL_TREE;
+
+  if ((TREE_CODE (t) == INDIRECT_REF
+       || TREE_CODE (t) == MEM_REF)
+      && TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == POINTER_TYPE)
+    {
+      attach_pt = TREE_OPERAND (t, 0);
+      if (TREE_CODE (attach_pt) == POINTER_PLUS_EXPR)
+	attach_pt = TREE_OPERAND (attach_pt, 0);
+    }
+
+  return attach_pt;
+}
+
 static const struct c_omp_directive omp_directives[] = {
   /* Keep this alphabetically sorted by the first word.  Non-null second/third
      if any should precede null ones.  */
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 4d156f6d3ec..cfac7d0a2b5 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13799,9 +13799,15 @@  handle_omp_array_sections (tree c, enum c_omp_region_type ort)
       if (size)
 	size = c_fully_fold (size, false, NULL);
       OMP_CLAUSE_SIZE (c) = size;
+      tree virtbase = t;
+      tree attach_pt
+	= ((ort != C_ORT_ACC)
+	   ? c_omp_decompose_attachable_address (t, &virtbase)
+	   : NULL_TREE);
       if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
 	  || (TREE_CODE (t) == COMPONENT_REF
-	      && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
+	      && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE
+	      && !attach_pt))
 	return false;
       gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR);
       switch (OMP_CLAUSE_MAP_KIND (c))
@@ -13834,10 +13840,10 @@  handle_omp_array_sections (tree c, enum c_omp_region_type ort)
       if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
 	  && !c_mark_addressable (t))
 	return false;
-      OMP_CLAUSE_DECL (c2) = t;
+      OMP_CLAUSE_DECL (c2) = attach_pt ? attach_pt : t;
       t = build_fold_addr_expr (first);
       t = fold_convert_loc (OMP_CLAUSE_LOCATION (c), ptrdiff_type_node, t);
-      tree ptr = OMP_CLAUSE_DECL (c2);
+      tree ptr = virtbase;
       if (!POINTER_TYPE_P (TREE_TYPE (ptr)))
 	ptr = build_fold_addr_expr (ptr);
       t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index e882c302f31..068c0c69e58 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5620,9 +5620,16 @@  handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	  OMP_CLAUSE_SIZE (c) = size;
 	  if (TREE_CODE (t) == FIELD_DECL)
 	    t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
+
+	  tree virtbase = t;
+	  tree attach_pt
+	    = ((ort != C_ORT_ACC)
+	       ? c_omp_decompose_attachable_address (t, &virtbase)
+	       : NULL_TREE);
 	  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
 	      || (TREE_CODE (t) == COMPONENT_REF
-		  && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
+		  && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE
+		  && !attach_pt))
 	    return false;
 	  switch (OMP_CLAUSE_MAP_KIND (c))
 	    {
@@ -5684,12 +5691,11 @@  handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	  if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
 	      && !cxx_mark_addressable (t))
 	    return false;
-	  OMP_CLAUSE_DECL (c2) = t;
+	  OMP_CLAUSE_DECL (c2) = attach_pt ? attach_pt : t;
 	  t = build_fold_addr_expr (first);
 	  t = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
 				ptrdiff_type_node, t);
-	  tree ptr = OMP_CLAUSE_DECL (c2);
-	  ptr = convert_from_reference (ptr);
+	  tree ptr = convert_from_reference (virtbase);
 	  if (!INDIRECT_TYPE_P (TREE_TYPE (ptr)))
 	    ptr = build_fold_addr_expr (ptr);
 	  t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
diff --git a/libgomp/testsuite/libgomp.c++/baseptrs-3.C b/libgomp/testsuite/libgomp.c++/baseptrs-3.C
new file mode 100644
index 00000000000..cabeb7c2b7a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/baseptrs-3.C
@@ -0,0 +1,182 @@ 
+/* { dg-xfail-if "fails to parse correctly" { *-*-* } } */
+
+#include <cstdlib>
+#include <cstring>
+#include <cassert>
+
+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 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->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);
+
+/* FIXME: This currently ICEs.  */
+/*  #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);
+
+  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->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);
+
+/* FIXME: This currently ICEs.  */
+/*  #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);
+
+  delete my_a->ptr;
+  delete my_a;
+  delete my_b;
+}
+
+int main (int argc, char *argv[])
+{
+  foo ();
+  bar ();
+  foop ();
+  barp ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/baseptrs-1.c b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-1.c
new file mode 100644
index 00000000000..073615625b7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-1.c
@@ -0,0 +1,50 @@ 
+#include <stdlib.h>
+#include <string.h>
+#include <assert.h>
+#include <stdio.h>
+
+#define N 32
+
+typedef struct {
+  int x2[10][N];
+} x1type;
+
+typedef struct {
+  x1type x1[10];
+} p2type;
+
+typedef struct {
+  p2type *p2;
+} p1type;
+
+typedef struct {
+  p1type *p1;
+} x0type;
+
+typedef struct {
+  x0type x0[10];
+} p0type;
+
+int main(int argc, char *argv[])
+{
+  p0type *p0;
+  int k1 = 0, k2 = 0, k3 = 0, n = N;
+
+  p0 = (p0type *) malloc (sizeof *p0);
+  p0->x0[0].p1 = (p1type *) malloc (sizeof *p0->x0[0].p1);
+  p0->x0[0].p1->p2 = (p2type *) malloc (sizeof *p0->x0[0].p1->p2);
+  memset (p0->x0[0].p1->p2, 0, sizeof *p0->x0[0].p1->p2);
+
+#pragma omp target map(tofrom: p0->x0[k1].p1->p2[k2].x1[k3].x2[4][0:n]) \
+		   map(to: p0->x0[k1].p1, p0->x0[k1].p1->p2) \
+		   map(to: p0->x0[k1].p1[0])
+  {
+    for (int i = 0; i < n; i++)
+      p0->x0[k1].p1->p2[k2].x1[k3].x2[4][i] = i;
+  }
+
+  for (int i = 0; i < n; i++)
+    assert (i == p0->x0[k1].p1->p2[k2].x1[k3].x2[4][i]);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/baseptrs-2.c b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-2.c
new file mode 100644
index 00000000000..e335d7da966
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-2.c
@@ -0,0 +1,70 @@ 
+#include <stdlib.h>
+#include <string.h>
+#include <assert.h>
+
+#define N 32
+
+typedef struct {
+  int arr[N];
+  int *ptr;
+} sc;
+
+typedef struct {
+  sc *c;
+} sb;
+
+typedef struct {
+  sb *b;
+  sc *c;
+} sa;
+
+int main (int argc, char *argv[])
+{
+  sa *p;
+
+  p = (sa *) malloc (sizeof *p);
+  p->b = (sb *) malloc (sizeof *p->b);
+  p->b->c = (sc *) malloc (sizeof *p->b->c);
+  p->c = (sc *) malloc (sizeof *p->c);
+  p->b->c->ptr = (int *) malloc (N * sizeof (int));
+  p->c->ptr = (int *) malloc (N * sizeof (int));
+
+  for (int i = 0; i < N; i++)
+    {
+      p->b->c->ptr[i] = 0;
+      p->c->ptr[i] = 0;
+      p->b->c->arr[i] = 0;
+      p->c->arr[i] = 0;
+    }
+
+#pragma omp target map(to: p->b, p->b[0], p->c, p->c[0], p->b->c, p->b->c[0]) \
+		   map(to: p->b->c->ptr, p->c->ptr) \
+		   map(tofrom: p->b->c->ptr[:N], p->c->ptr[:N])
+  {
+    for (int i = 0; i < N; i++)
+      {
+	p->b->c->ptr[i] = i;
+	p->c->ptr[i] = i * 2;
+      }
+  }
+
+#pragma omp target map(to: p->b, p->b[0], p->b->c, p->c) \
+		   map(tofrom: p->c[0], p->b->c[0])
+  {
+    for (int i = 0; i < N; i++)
+      {
+	p->b->c->arr[i] = i * 3;
+	p->c->arr[i] = i * 4;
+      }
+  }
+
+  for (int i = 0; i < N; i++)
+    {
+      assert (p->b->c->ptr[i] == i);
+      assert (p->c->ptr[i] == i * 2);
+      assert (p->b->c->arr[i] == i * 3);
+      assert (p->c->arr[i] == i * 4);
+    }
+
+  return 0;
+}