[PATCH, OpenMP, C/C++] Handle array reference base-pointers in array sections

[Chung-Lin Tang <cltang@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 1721 bytes --]

Hi Jakub,
as encountered in cases where a program constructs its own deep-copying
for arrays-of-pointers, e.g:

    #pragma omp target enter data map(to:level->vectors[:N])
    for (i = 0; i < N; i++)
      #pragma omp target enter data map(to:level->vectors[i][:N])

We need to treat the part of the array reference before the array section
as a base-pointer (here 'level->vectors[i]'), providing pointer-attachment behavior.

This patch adds this inside handle_omp_array_sections(), tracing the whole
sequence of array dimensions, creating a whole base-pointer reference
iteratively using build_array_ref(). The conditions are that each of the
"absorbed" dimensions must be length==1, and the final reference must be
of pointer-type (so that pointer attachment makes sense).

There's also a little patch in gimplify_scan_omp_clauses(), to make sure
the array-ref base-pointer goes down the right path.

This case was encountered when working to make 534.hpgmgfv_t from
SPEChpc 2021 properly compile. Tested without regressions on trunk.
Okay to go in once stage1 opens?

Thanks,
Chung-Lin

2022-02-21  Chung-Lin Tang  <cltang@codesourcery.com>

gcc/c/ChangeLog:

	* c-typeck.cc (handle_omp_array_sections): Add handling for
	creating array-reference base-pointer attachment clause.

gcc/cp/ChangeLog:

	* semantics.cc (handle_omp_array_sections): Add handling for
	creating array-reference base-pointer attachment clause.

gcc/ChangeLog:

	* gimplify.cc (gimplify_scan_omp_clauses): Add case for
	attach/detach map kind for ARRAY_REF of POINTER_TYPE.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/target-enter-data-1.c: Adjust testcase.

libgomp/testsuite/ChangeLog:

	* libgomp.c-c++-common/ptr-attach-2.c: New test.

[-- Attachment #2: arrayref-bptr.patch --]
[-- Type: text/plain, Size: 6794 bytes --]

diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 3075c883548..4257e373557 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -13649,6 +13649,10 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
       if (int_size_in_bytes (TREE_TYPE (first)) <= 0)
 	maybe_zero_len = true;
 
+      struct dim { tree low_bound, length; };
+      auto_vec<dim> dims (num);
+      dims.safe_grow (num);
+
       for (i = num, t = OMP_CLAUSE_DECL (c); i > 0;
 	   t = TREE_CHAIN (t))
 	{
@@ -13763,6 +13767,9 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	      else
 		size = size_binop (MULT_EXPR, size, l);
 	    }
+
+	  dim d = { low_bound, length };
+	  dims[i] = d;
 	}
       if (side_effects)
 	size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
@@ -13802,6 +13809,23 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	  OMP_CLAUSE_DECL (c) = t;
 	  return false;
 	}
+
+      tree aref = t;
+      for (i = 0; i < dims.length (); i++)
+	{
+	  if (dims[i].length && integer_onep (dims[i].length))
+	    {
+	      tree lb = dims[i].low_bound;
+	      aref = build_array_ref (OMP_CLAUSE_LOCATION (c), aref, lb);
+	    }
+	  else
+	    {
+	      if (TREE_CODE (TREE_TYPE (aref)) == POINTER_TYPE)
+		t = aref;
+	      break;
+	    }
+	}
+
       first = c_fully_fold (first, false, NULL);
       OMP_CLAUSE_DECL (c) = first;
       if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
@@ -13836,7 +13860,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	  break;
 	}
       tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
-      if (TREE_CODE (t) == COMPONENT_REF)
+      if (TREE_CODE (t) == COMPONENT_REF || TREE_CODE (t) == ARRAY_REF
+	  || TREE_CODE (t) == INDIRECT_REF)
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
       else
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 0cb17a6a8ab..646f4883d66 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -5497,6 +5497,10 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
       if (processing_template_decl && maybe_zero_len)
 	return false;
 
+      struct dim { tree low_bound, length; };
+      auto_vec<dim> dims (num);
+      dims.safe_grow (num);
+
       for (i = num, t = OMP_CLAUSE_DECL (c); i > 0;
 	   t = TREE_CHAIN (t))
 	{
@@ -5604,6 +5608,9 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	      else
 		size = size_binop (MULT_EXPR, size, l);
 	    }
+
+	  dim d = { low_bound, length };
+	  dims[i] = d;
 	}
       if (!processing_template_decl)
 	{
@@ -5647,6 +5654,24 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	      OMP_CLAUSE_DECL (c) = t;
 	      return false;
 	    }
+
+	  tree aref = t;
+	  for (i = 0; i < dims.length (); i++)
+	    {
+	      if (dims[i].length && integer_onep (dims[i].length))
+		{
+		  tree lb = dims[i].low_bound;
+		  aref = convert_from_reference (aref);
+		  aref = build_array_ref (OMP_CLAUSE_LOCATION (c), aref, lb);
+		}
+	      else
+		{
+		  if (TREE_CODE (TREE_TYPE (aref)) == POINTER_TYPE)
+		    t = aref;
+		  break;
+		}
+	    }
+
 	  OMP_CLAUSE_DECL (c) = first;
 	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
 	    return false;
@@ -5681,7 +5706,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	  bool reference_always_pointer = true;
 	  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
 				      OMP_CLAUSE_MAP);
-	  if (TREE_CODE (t) == COMPONENT_REF)
+	  if (TREE_CODE (t) == COMPONENT_REF || TREE_CODE (t) == ARRAY_REF
+	      || (TREE_CODE (t) == INDIRECT_REF && !REFERENCE_REF_P (t)))
 	    {
 	      OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
 
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index f570daa015a..77b95cd8000 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9626,7 +9626,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		   || (component_ref_p
 		       && (INDIRECT_REF_P (decl)
 			   || TREE_CODE (decl) == MEM_REF
-			   || TREE_CODE (decl) == ARRAY_REF)))
+			   || TREE_CODE (decl) == ARRAY_REF))
+		   || (TREE_CODE (decl) == ARRAY_REF
+		       && TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE
+		       && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH))
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
diff --git a/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c b/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c
index ce766d29e2d..3a1b488fa1f 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c
@@ -21,4 +21,5 @@ void func (struct foo *f, int n, int m)
   #pragma omp target enter data map (to: f->bars[n].vectors[:f->bars[n].num_vectors])
 }
 
-/* { dg-final { scan-tree-dump-times "map\\(to:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:\[^-\]+->vectors \\\[bias: \[^\]\]+\\\]\\)" 3 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(to:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:\\*_\[0-9\]+ \\\[bias: \[^\]\]+\\\]\\)" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(to:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:\[^-\]+->vectors \\\[bias: \[^\]\]+\\\]\\)" 2 "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-2.c b/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-2.c
new file mode 100644
index 00000000000..889a4a253ae
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/ptr-attach-2.c
@@ -0,0 +1,60 @@
+#include <stdlib.h>
+
+struct blk { int x, y; };
+struct L
+{
+  #define N 10
+  struct {
+    int num_blocks[N];
+    struct blk * blocks[N];
+  } m;
+};
+
+void foo (struct L *l)
+{
+  for (int i = 0; i < N; i++)
+    {
+      l->m.blocks[i] = (struct blk *) malloc (sizeof (struct blk) * N);
+      l->m.num_blocks[i] = N;
+    }
+
+  #pragma omp target enter data map(to:l[:1])
+  for (int i = 0; i < N; i++)
+    {
+      #pragma omp target enter data map(to:l->m.blocks[i][:l->m.num_blocks[i]])
+    }
+
+  #pragma omp target
+  {
+    for (int i = 0; i < N; i++)
+      for (int j = 0; j < N; j++)
+	{
+	  l->m.blocks[i][j].x = i + j;
+	  l->m.blocks[i][j].y = i * j;
+	}
+  }
+
+  for (int i = 0; i < N; i++)
+    {
+      #pragma omp target exit data map(from:l->m.blocks[i][:l->m.num_blocks[i]])
+    }
+  #pragma omp target exit data map(from:l[:1])
+
+
+  for (int i = 0; i < N; i++)
+    for (int j = 0; j < N; j++)
+      {
+	if (l->m.blocks[i][j].x != i + j)
+	  abort ();
+	if (l->m.blocks[i][j].y != i * j)
+	  abort ();
+      }
+
+}
+
+int main (void)
+{
+  struct L l;
+  foo (&l);
+  return 0;
+}