OpenACC: whole struct vs. component mappings (PR107028)

This patch fixes an ICE when both a complete struct variable and
components of that struct are mapped on the same directive for OpenACC,
using a modified version of the scheme used for OpenMP in the following
patch:

  https://gcc.gnu.org/pipermail/gcc-patches/2022-September/601558.html

A new function has been added to make sure that the mapping kinds of
the whole struct and the member access are compatible -- conservatively,
so as not to copy more to/from the device than the user expects.

This version of the patch uses a different method to detect duplicate
clauses for OpenACC in oacc_resolve_clause_dependencies, and removes
the now-redundant check in omp_accumulate_sibling_lists.  (The latter
check would no longer trigger when we map the whole struct on the same
directive because the component-mapping clauses are now deleted before
the check is executed.)

2022-09-28  Julian Brown  <julian@codesourcery.com>

gcc/
	PR middle-end/107028
	* gimplify.cc (omp_check_mapping_compatibility,
	oacc_resolve_clause_dependencies): New functions.
	(omp_accumulate_sibling_list): Remove redundant duplicate clause
	detection for OpenACC.
	(build_struct_sibling_lists): Skip deleted groups.  Don't build sibling
	list for struct variables that are fully mapped on the same directive
	for OpenACC.
	(gimplify_scan_omp_clauses): Call oacc_resolve_clause_dependencies.

gcc/testsuite/
	PR middle-end/107028
	* c-c++-common/goacc/struct-component-kind-1.c: New test.
	* g++.dg/goacc/pr107028-1.C: New test.
	* g++.dg/goacc/pr107028-2.C: New test.
	* gfortran.dg/goacc/mapping-tests-5.f90: New test.

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 4d032c6bf06372d07b27dfae740776417087ef79..90606088128985e4469ec1c22a157d299e1ff260 100644 (file)
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9861,6 +9861,134 @@ omp_lastprivate_for_combined_outer_constructs (struct gimplify_omp_ctx *octx,
     omp_notice_variable (octx, decl, true);
 }
 
+/* If we have mappings INNER and OUTER, where INNER is a component access and
+   OUTER is a mapping of the whole containing struct, check that the mappings
+   are compatible.  We'll be deleting the inner mapping, so we need to make
+   sure the outer mapping does (at least) the same transfers to/from the device
+   as the inner mapping.  */
+
+bool
+omp_check_mapping_compatibility (location_t loc,
+                                omp_mapping_group *outer,
+                                omp_mapping_group *inner)
+{
+  tree first_outer = *outer->grp_start, first_inner = *inner->grp_start;
+
+  gcc_assert (OMP_CLAUSE_CODE (first_outer) == OMP_CLAUSE_MAP);
+  gcc_assert (OMP_CLAUSE_CODE (first_inner) == OMP_CLAUSE_MAP);
+
+  enum gomp_map_kind outer_kind = OMP_CLAUSE_MAP_KIND (first_outer);
+  enum gomp_map_kind inner_kind = OMP_CLAUSE_MAP_KIND (first_inner);
+
+  if (outer_kind == inner_kind)
+    return true;
+
+  switch (outer_kind)
+    {
+    case GOMP_MAP_ALWAYS_TO:
+      if (inner_kind == GOMP_MAP_FORCE_PRESENT
+         || inner_kind == GOMP_MAP_ALLOC
+         || inner_kind == GOMP_MAP_TO)
+       return true;
+      break;
+
+    case GOMP_MAP_ALWAYS_FROM:
+      if (inner_kind == GOMP_MAP_FORCE_PRESENT
+         || inner_kind == GOMP_MAP_ALLOC
+         || inner_kind == GOMP_MAP_FROM)
+       return true;
+      break;
+
+    case GOMP_MAP_TO:
+    case GOMP_MAP_FROM:
+      if (inner_kind == GOMP_MAP_FORCE_PRESENT
+         || inner_kind == GOMP_MAP_ALLOC)
+       return true;
+      break;
+
+    case GOMP_MAP_ALWAYS_TOFROM:
+    case GOMP_MAP_TOFROM:
+      if (inner_kind == GOMP_MAP_FORCE_PRESENT
+         || inner_kind == GOMP_MAP_ALLOC
+         || inner_kind == GOMP_MAP_TO
+         || inner_kind == GOMP_MAP_FROM
+         || inner_kind == GOMP_MAP_TOFROM)
+       return true;
+      break;
+
+    default:
+      ;
+    }
+
+  error_at (loc, "data movement for component %qE is not compatible with "
+           "movement for struct %qE", OMP_CLAUSE_DECL (first_inner),
+           OMP_CLAUSE_DECL (first_outer));
+
+  return false;
+}
+
+/* Similar to omp_resolve_clause_dependencies, but for OpenACC.  The only
+   clause dependencies we handle for now are struct element mappings and
+   whole-struct mappings on the same directive, and duplicate clause
+   detection.  */
+
+void
+oacc_resolve_clause_dependencies (vec<omp_mapping_group> *groups,
+                                 hash_map<tree_operand_hash,
+                                          omp_mapping_group *> *grpmap)
+{
+  int i;
+  omp_mapping_group *grp;
+  hash_set<tree_operand_hash> *seen_components = NULL;
+  hash_set<tree_operand_hash> *shown_error = NULL;
+
+  FOR_EACH_VEC_ELT (*groups, i, grp)
+    {
+      tree grp_end = grp->grp_end;
+      tree decl = OMP_CLAUSE_DECL (grp_end);
+
+      gcc_assert (OMP_CLAUSE_CODE (grp_end) == OMP_CLAUSE_MAP);
+
+      if (DECL_P (grp_end))
+       continue;
+
+      tree c = OMP_CLAUSE_DECL (*grp->grp_start);
+      while (TREE_CODE (c) == ARRAY_REF)
+       c = TREE_OPERAND (c, 0);
+      if (TREE_CODE (c) != COMPONENT_REF)
+       continue;
+      if (!seen_components)
+       seen_components = new hash_set<tree_operand_hash> ();
+      if (!shown_error)
+       shown_error = new hash_set<tree_operand_hash> ();
+      if (seen_components->contains (c)
+         && !shown_error->contains (c))
+       {
+         error_at (OMP_CLAUSE_LOCATION (grp_end),
+                   "%qE appears more than once in map clauses",
+                   OMP_CLAUSE_DECL (grp_end));
+         shown_error->add (c);
+       }
+      else
+       seen_components->add (c);
+
+      omp_mapping_group *struct_group;
+      if (omp_mapped_by_containing_struct (grpmap, decl, &struct_group)
+         && *grp->grp_start == grp_end)
+       {
+         omp_check_mapping_compatibility (OMP_CLAUSE_LOCATION (grp_end),
+                                          struct_group, grp);
+         /* Remove the whole of this mapping -- redundant.  */
+         grp->deleted = true;
+       }
+    }
+
+  if (seen_components)
+    delete seen_components;
+  if (shown_error)
+    delete shown_error;
+}
+
 /* Link node NEWNODE so it is pointed to by chain INSERT_AT.  NEWNODE's chain
    is linked to the previous node pointed to by INSERT_AT.  */
 
@@ -10238,37 +10366,6 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
              break;
            if (scp)
              continue;
-           if ((region_type & ORT_ACC) != 0)
-             {
-               /* This duplicate checking code is currently only enabled for
-                  OpenACC.  */
-               tree d1 = OMP_CLAUSE_DECL (*sc);
-               tree d2 = OMP_CLAUSE_DECL (grp_end);
-               while (TREE_CODE (d1) == ARRAY_REF)
-                 d1 = TREE_OPERAND (d1, 0);
-               while (TREE_CODE (d2) == ARRAY_REF)
-                 d2 = TREE_OPERAND (d2, 0);
-               if (TREE_CODE (d1) == INDIRECT_REF)
-                 d1 = TREE_OPERAND (d1, 0);
-               if (TREE_CODE (d2) == INDIRECT_REF)
-                 d2 = TREE_OPERAND (d2, 0);
-               while (TREE_CODE (d1) == COMPONENT_REF)
-                 if (TREE_CODE (d2) == COMPONENT_REF
-                     && TREE_OPERAND (d1, 1) == TREE_OPERAND (d2, 1))
-                   {
-                     d1 = TREE_OPERAND (d1, 0);
-                     d2 = TREE_OPERAND (d2, 0);
-                   }
-                 else
-                   break;
-               if (d1 == d2)
-                 {
-                   error_at (OMP_CLAUSE_LOCATION (grp_end),
-                             "%qE appears more than once in map clauses",
-                             OMP_CLAUSE_DECL (grp_end));
-                   return NULL;
-                 }
-             }
            if (maybe_lt (coffset, offset)
                || (known_eq (coffset, offset)
                    && maybe_lt (cbitpos, bitpos)))
@@ -10400,6 +10497,11 @@ omp_build_struct_sibling_lists (enum tree_code code,
       if (DECL_P (decl))
        continue;
 
+      /* Skip groups we marked for deletion in
+        oacc_resolve_clause_dependencies.  */
+      if (grp->deleted)
+       continue;
+
       if (OMP_CLAUSE_CHAIN (*grp_start_p)
          && OMP_CLAUSE_CHAIN (*grp_start_p) != grp_end)
        {
@@ -10436,14 +10538,14 @@ omp_build_struct_sibling_lists (enum tree_code code,
       if (TREE_CODE (decl) != COMPONENT_REF)
        continue;
 
-      /* If we're mapping the whole struct in another node, skip creation of
-        sibling lists.  */
+      /* If we're mapping the whole struct in another node, skip adding this
+        node to a sibling list.  */
       omp_mapping_group *wholestruct;
-      if (!(region_type & ORT_ACC)
-         && omp_mapped_by_containing_struct (*grpmap, OMP_CLAUSE_DECL (c),
-                                             &wholestruct))
+      if (omp_mapped_by_containing_struct (*grpmap, OMP_CLAUSE_DECL (c),
+                                          &wholestruct))
        {
-         if (*grp_start_p == grp_end)
+         if (!(region_type & ORT_ACC)
+             && *grp_start_p == grp_end)
            /* Remove the whole of this mapping -- redundant.  */
            grp->deleted = true;
 
@@ -10632,6 +10734,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
          hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
          grpmap = omp_index_mapping_groups (groups);
 
+         oacc_resolve_clause_dependencies (groups, grpmap);
          omp_build_struct_sibling_lists (code, region_type, groups, &grpmap,
                                          list_p);
 

diff --git a/gcc/testsuite/c-c++-common/goacc/struct-component-kind-1.c b/gcc/testsuite/c-c++-common/goacc/struct-component-kind-1.c
new file mode 100644 (file)
index 0000000..8d2f5ea
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/struct-component-kind-1.c
@@ -0,0 +1,72 @@
+/* { dg-do compile } */
+
+#include <stdlib.h>
+
+#define N 20
+
+struct s {
+  int base[N];
+};
+
+int main (void)
+{
+  struct s v;
+
+#pragma acc parallel copy(v, v.base[0:N])
+{ }
+
+#pragma acc parallel copyin(v, v.base[0:N])
+{ }
+
+#pragma acc parallel copyout(v, v.base[0:N])
+{ }
+
+#pragma acc parallel copy(v) copyin(v.base[0:N])
+{ }
+
+#pragma acc parallel copy(v) copyout(v.base[0:N])
+{ }
+
+#pragma acc parallel copy(v) present(v.base[0:N])
+{ }
+
+#pragma acc parallel copyin(v) present(v.base[0:N])
+{ }
+
+#pragma acc parallel copyout(v) present(v.base[0:N])
+{ }
+
+#pragma acc enter data copyin(v, v.base[0:N])
+#pragma acc update device(v, v.base[0:N])
+#pragma acc exit data delete(v, v.base[0:N])
+
+#pragma acc parallel copyin(v) copy(v.base[0:N])
+/* { dg-error "data movement for component 'v\\.(s::)?base\\\[0\\\]' is not compatible with movement for struct 'v'" "" { target *-*-* } .-1 } */
+{ }
+
+#pragma acc parallel copyout(v) copy(v.base[0:N])
+/* { dg-error "data movement for component 'v\\.(s::)?base\\\[0\\\]' is not compatible with movement for struct 'v'" "" { target *-*-* } .-1 } */
+{ }
+
+#pragma acc parallel present(v) copy(v.base[0:N])
+/* { dg-error "data movement for component 'v\\.(s::)?base\\\[0\\\]' is not compatible with movement for struct 'v'" "" { target *-*-* } .-1 } */
+{ }
+
+#pragma acc parallel present(v) copyin(v.base[0:N])
+/* { dg-error "data movement for component 'v\\.(s::)?base\\\[0\\\]' is not compatible with movement for struct 'v'" "" { target *-*-* } .-1 } */
+{ }
+
+#pragma acc parallel present(v) copyout(v.base[0:N])
+/* { dg-error "data movement for component 'v\\.(s::)?base\\\[0\\\]' is not compatible with movement for struct 'v'" "" { target *-*-* } .-1 } */
+{ }
+
+#pragma acc parallel present(v) no_create(v.base[0:N])
+/* { dg-error "data movement for component 'v\\.(s::)?base\\\[0\\\]' is not compatible with movement for struct 'v'" "" { target *-*-* } .-1 } */
+{ }
+
+#pragma acc parallel no_create(v) present(v.base[0:N])
+/* { dg-error "data movement for component 'v\\.(s::)?base\\\[0\\\]' is not compatible with movement for struct 'v'" "" { target *-*-* } .-1 } */
+{ }
+
+  return 0;
+}

diff --git a/gcc/testsuite/g++.dg/goacc/pr107028-1.C b/gcc/testsuite/g++.dg/goacc/pr107028-1.C
new file mode 100644 (file)
index 0000000..93b8743
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/pr107028-1.C
@@ -0,0 +1,14 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+
+class data_container {
+ public:
+  int data;
+};
+
+void test2() {
+  data_container a;
+#pragma acc data copyin(a, a.data)
+// { dg-final { scan-tree-dump {map\(to:a \[len: [0-9]+\]\)} "gimple" } }
+{ }
+}

diff --git a/gcc/testsuite/g++.dg/goacc/pr107028-2.C b/gcc/testsuite/g++.dg/goacc/pr107028-2.C
new file mode 100644 (file)
index 0000000..cf741bd
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/pr107028-2.C
@@ -0,0 +1,27 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+
+#include <cstdlib>
+
+typedef float real_t;
+
+struct foo {
+  real_t *data;
+};
+
+#define n 1024
+
+int test3() {
+    real_t *a = (real_t *)malloc(n * sizeof(real_t));
+    struct foo b;
+    b.data = (real_t *)malloc(n * sizeof(real_t));
+
+    #pragma acc data copyin(a[0:n], b, b.data[0:n])
+// { dg-final { scan-tree-dump {map\(to:\*_[0-9]+ \[len: [0-9]+\]\) map\(attach:b.data \[bias: 0\]\) map\(to:b \[len: [0-9]+\]\) map\(to:\*a \[len: [0-9]+\]\)} "gimple" } }
+    { }
+
+    free (b.data);
+    free (a);
+
+    return 0;
+}

diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-5.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-5.f90
new file mode 100644 (file)
index 0000000..8df8c58
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-5.f90
@@ -0,0 +1,15 @@
+subroutine foo
+  type one
+    integer, dimension(10) :: i, j
+  end type
+  type two
+    type(one) A, B
+  end type
+
+  type(two) x
+
+  !$acc enter data copyin(x%A%i(5), x%A%i(4), x%A)
+! { dg-error ".x.a.i. appears more than once in map clauses" "" { target *-*-* } .-1 }
+  !$acc enter data copyin(x%A, x%A%i(5), x%A%i(4))
+! { dg-error ".x.a.i. appears more than once in map clauses" "" { target *-*-* } .-1 }
+end