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/ChangeLog:

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

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index d093c79efc66803f783db37f5651a2876d036e4e..1bfa4f6a664f387ddb50b464270616fda3163eb3 100644 (file)
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,8 @@
+2022-02-24  Chung-Lin Tang  <cltang@codesourcery.com>
+
+       * gimplify.cc (gimplify_scan_omp_clauses): Add case for
+       attach/detach map kind for ARRAY_REF of POINTER_TYPE.
+
 2022-02-11  Kwok Cheung Yeung  <kcy@codesourcery.com>
 
        * omp-general.cc (DELAY_METADIRECTIVES_AFTER_LTO): Check that cfun is

diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp
index 3e5fe5fd85e8438d94450828d94b0d999294174b..11d0bbe781813a54d33f8adf630a71f57948d97a 100644 (file)
--- a/gcc/c/ChangeLog.omp
+++ b/gcc/c/ChangeLog.omp
@@ -1,3 +1,8 @@
+2022-02-24  Chung-Lin Tang  <cltang@codesourcery.com>
+
+       * c-typeck.cc (handle_omp_array_sections): Add handling for
+       creating array-reference base-pointer attachment clause.
+
 2022-02-18  Kwok Cheung Yeung  <kcy@codesourcery.com>
 
        * c-parser.cc (c_parser_omp_construct): Move handling of

diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 32bf33e80d6955f4668d9870e26495fd50713774..c974a6095133c745e46bddd55deb85bada3bf240 100644 (file)
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -13677,6 +13677,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))
        {
@@ -13798,6 +13802,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)
        {
@@ -13845,6 +13852,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)
@@ -13879,7 +13903,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/ChangeLog.omp b/gcc/cp/ChangeLog.omp
index 251fc73f190a0187280cd9451fdf5f1ad32293a3..b0f96782d48e7eecea7e3bae251b9f7807d2785d 100644 (file)
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,8 @@
+2022-02-24  Chung-Lin Tang  <cltang@codesourcery.com>
+
+       * semantics.cc (handle_omp_array_sections): Add handling for
+       creating array-reference base-pointer attachment clause.
+
 2022-02-18  Kwok Cheung Yeung  <kcy@codesourcery.com>
 
        * parser.cc (cp_parser_omp_construct): Move handling of

diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 539bdca7eb0fee814110f03c50ee4025e5f6ba3e..49cdaf78d731b87155054e4de2bfba622048fc28 100644 (file)
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -5526,6 +5526,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))
        {
@@ -5645,6 +5649,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)
        {
@@ -5696,6 +5703,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;
@@ -5730,7 +5755,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 105bd81b5b8b3d558e9d2c71e515de77dd186246..543b9445092aca59099840251eb13ee4b7b6132c 100644 (file)
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -10490,7 +10490,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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 72c03338b0202b7850cfbebc479431da739159b4..3b9e248f5fa35f6907c9a99431f8bbc9d531a93d 100644 (file)
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,7 @@
+2022-02-24  Chung-Lin Tang  <cltang@codesourcery.com>
+
+       * c-c++-common/gomp/target-enter-data-1.c: Adjust testcase.
+
 2022-02-18  Kwok Cheung Yeung  <kcy@codesourcery.com>
 
        * c-c++-common/gomp/metadirective-1.c (f): Add test for

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/ChangeLog.omp b/libgomp/ChangeLog.omp
index 5f4fb9362b596a62ed635abd3188f5756026206c..865d7a839948eb79dfd20f350938da5e95606c46 100644 (file)
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,7 @@
+2022-02-24  Chung-Lin Tang  <cltang@codesourcery.com>
+
+       * testsuite/libgomp.c-c++-common/ptr-attach-2.c: New test.
+
 2022-01-28  Kwok Cheung Yeung  <kcy@codesourcery.com>
 
        * testsuite/libgomp.c-c++-common/metadirective-2.c (main): Add

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;
+}