openmp: Handle C/C++ array reference base-pointers in array sections

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).

Merged from:
https://gcc.gnu.org/pipermail/gcc-patches/2022-February/590658.html

2022-02-24  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.

diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 29f53a926e334c153b198995e1ca53fdbc67bcfd..7b40f5ae8dfcd57f8b47ac6aacf793174829a544 100644 (file)
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13547,6 +13547,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))
        {
@@ -13668,6 +13672,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 (non_contiguous)
        {
@@ -13715,6 +13722,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 (size)
@@ -13747,7 +13771,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.c b/gcc/cp/semantics.c
index 59c052d8eb638c614df343f0d9bd7afff53583d3..87f996fec1f1acb03986ae98d0bddaf6a8619990 100644 (file)
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5419,6 +5419,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))
        {
@@ -5538,6 +5542,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)
        {
@@ -5589,6 +5596,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;
          OMP_CLAUSE_SIZE (c) = size;
          if (TREE_CODE (t) == FIELD_DECL)
@@ -5621,7 +5646,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.c b/gcc/gimplify.c
index 6e0c900e26c9526addbf79f94707e4f12cc286bd..b0226e98d02c16116ce4db7a77aeb929930add31 100644 (file)
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10224,7 +10224,11 @@ 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 ce766d29e2dc245cffde13c5a771e14d6c365c0e..3a1b488fa1f8e7f045201c33f0745a31b5a67126 100644 (file)
--- 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 (file)
index 0000000..889a4a2
--- /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;
+}