OpenACC: whole struct vs. component mappings (PR107028)

Message ID 20220928132024.64984-1-julian@codesourcery.com
State New
Headers
Series OpenACC: whole struct vs. component mappings (PR107028) |

Commit Message

Julian Brown Sept. 28, 2022, 1:20 p.m. UTC
  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.

Tested with offloading to NVPTX. OK?

Thanks,

Julian

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

gcc/
	PR middle-end/107028
	* gimplify.cc (omp_group_base): Fix IF_PRESENT handling.
	(omp_check_mapping_compatibility, oacc_resolve_clause_dependencies):
	New functions.
	(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.
---
 gcc/gimplify.cc                               | 129 +++++++++++++++++-
 .../goacc/struct-component-kind-1.c           |  72 ++++++++++
 gcc/testsuite/g++.dg/goacc/pr107028-1.C       |  14 ++
 gcc/testsuite/g++.dg/goacc/pr107028-2.C       |  27 ++++
 4 files changed, 235 insertions(+), 7 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/struct-component-kind-1.c
 create mode 100644 gcc/testsuite/g++.dg/goacc/pr107028-1.C
 create mode 100644 gcc/testsuite/g++.dg/goacc/pr107028-2.C
  

Comments

Tobias Burnus Sept. 28, 2022, 3:17 p.m. UTC | #1
On 28.09.22 15:20, Julian Brown wrote:

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 [...]
Tested with offloading to NVPTX. OK?

OpenACC comments:

I do note that there are now two "appears more than once in map clauses". The
newly added  error_at in oacc_resolve_clause_dependencies is triggered by
gcc/testsuite/gfortran.dg/goacc/{derived-types-3.f90,goacc/mapping-tests-{1,4}.f90}.

However, it looks as if no testcase triggers anymore the OpenACC-only error_at
in omp_accumulate_sibling_list. My impression is that the oacc_resolve_clause_dependencies
check + 'if (grp->deleted) continue;' in the sibling-list function makes this difficult to hit.

I don't see immediately whether some cases can still reach omp_accumulate_sibling_list –
if so, a testcase would be nice, or whether that error_at can now be removed.


However, I note that *without* the patch, the *following* *error* triggers – while
it compiles *silently* *with* the *patch* applied:

   15 |   !$acc enter data copyin(x%A, x%A%i(5), x%A%i(5))
      |                                                  ^
Error: ‘x.a.i’ appears more than once in map clauses

   15 |   !$acc enter data copyin(x%A, x%A%i(5), x%A%i(4))
      |                                                  ^
Error: ‘x.a.i’ appears more than once in map clauses

BTW: The two testcases differ by the array-element: '5'/'5' vs. '5'/'4'.
Testcase is a modified existing one:

--- a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
@@ -3 +3 @@ subroutine foo
-    integer i, j
+    integer i(5), j
@@ -15 +15 @@ subroutine foo
-  !$acc enter data copyin(x%A, x%A%i, x%A%i)
+  !$acc enter data copyin(x%A, x%A%i(5), x%A%i(4))

Otherwise, the patch looks sensible - as far as I understand it.
However, it surely would help if Thomas and/or Jakub could have a second look.

Tobias


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

gcc/
        PR middle-end/107028
        * gimplify.cc (omp_group_base): Fix IF_PRESENT handling.
        (omp_check_mapping_compatibility, oacc_resolve_clause_dependencies):
        New functions.
        (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.

-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
  
Julian Brown Sept. 29, 2022, 12:59 p.m. UTC | #2
On Wed, 28 Sep 2022 17:17:30 +0200
Tobias Burnus <tobias@codesourcery.com> wrote:

> On 28.09.22 15:20, Julian Brown wrote:
> 
> 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 [...]
> Tested with offloading to NVPTX. OK?
> 
> OpenACC comments:
> 
> I do note that there are now two "appears more than once in map
> clauses". The newly added  error_at in
> oacc_resolve_clause_dependencies is triggered by
> gcc/testsuite/gfortran.dg/goacc/{derived-types-3.f90,goacc/mapping-tests-{1,4}.f90}.

> I don't see immediately whether some cases can still reach
> omp_accumulate_sibling_list – if so, a testcase would be nice, or
> whether that error_at can now be removed.

This version of the patch removes the now-redundant check in
omp_accumulate_sibling_list.

> However, I note that *without* the patch, the *following* *error*
> triggers – while it compiles *silently* *with* the *patch* applied:
> 
>    15 |   !$acc enter data copyin(x%A, x%A%i(5), x%A%i(5))
>       |                                                  ^
> Error: ‘x.a.i’ appears more than once in map clauses
> 
>    15 |   !$acc enter data copyin(x%A, x%A%i(5), x%A%i(4))
>       |                                                  ^
> Error: ‘x.a.i’ appears more than once in map clauses
> 
> BTW: The two testcases differ by the array-element: '5'/'5' vs.
> '5'/'4'. Testcase is a modified existing one:

...and this test now triggers an error again (as it should -- you can't
map more than one part of the same array). Slightly unfortunately we're
not using the existing "group map" any more, since it doesn't record
quite the right thing -- instead, a local hash set is used to detect
duplicates in oacc_resolve_clause_dependencies.

Re-tested with offloading to NVPTX. OK?

Thanks,

Julian
  
Tobias Burnus Sept. 29, 2022, 1:09 p.m. UTC | #3
On 29.09.22 14:59, Julian Brown wrote:
> On Wed, 28 Sep 2022 17:17:30 +0200 Tobias Burnus <tobias@codesourcery.com> wrote:
>> I don't see immediately whether some cases can still reach
>> omp_accumulate_sibling_list – if so, a testcase would be nice, or
>> whether that error_at can now be removed.
> This version of the patch removes the now-redundant check in
> omp_accumulate_sibling_list.
Thanks!
>> However, I note that *without* the patch, the *following* *error*
>> triggers – while it compiles *silently* *with* the *patch* applied:
>> [...]
> ...and this test now triggers an error again (as it should -- you can't
> map more than one part of the same array).
...
> Re-tested with offloading to NVPTX. OK?

LGTM.

Thanks for the patch!

Tobias

-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
  

Patch

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 4d032c6bf06..7fc1d38644a 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9245,6 +9245,7 @@  omp_group_base (omp_mapping_group *grp, unsigned int *chained,
     case GOMP_MAP_RELEASE:
     case GOMP_MAP_DELETE:
     case GOMP_MAP_FORCE_ALLOC:
+    case GOMP_MAP_IF_PRESENT:
       if (node == grp->grp_end)
 	return node;
 
@@ -9323,7 +9324,6 @@  omp_group_base (omp_mapping_group *grp, unsigned int *chained,
     case GOMP_MAP_FORCE_DEVICEPTR:
     case GOMP_MAP_DEVICE_RESIDENT:
     case GOMP_MAP_LINK:
-    case GOMP_MAP_IF_PRESENT:
     case GOMP_MAP_FIRSTPRIVATE:
     case GOMP_MAP_FIRSTPRIVATE_INT:
     case GOMP_MAP_USE_DEVICE_PTR:
@@ -9861,6 +9861,115 @@  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 the above function, but for OpenACC.  The only clause
+   dependencies we handle for now are struct element mappings and whole-struct
+   mappings on the same directive.  */
+
+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;
+
+  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);
+
+      omp_mapping_group **maybe_siblings = grpmap->get (decl);
+
+      if (maybe_siblings
+	  && !(*maybe_siblings)->deleted
+	  && (*maybe_siblings)->sibling)
+	{
+	  error_at (OMP_CLAUSE_LOCATION (grp_end),
+		    "%qE appears more than once in map clauses",
+		    OMP_CLAUSE_DECL (grp_end));
+	  (*maybe_siblings)->deleted = true;
+	}
+
+      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;
+	}
+    }
+}
+
 /* 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.  */
 
@@ -10400,6 +10509,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 +10550,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 +10746,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
index 00000000000..8d2f5ea6497
--- /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
index 00000000000..93b87439b4f
--- /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
index 00000000000..cf741bd78c7
--- /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;
+}