* [PATCH] OpenACC: whole struct vs. component mappings (PR107028)
@ 2022-09-28 13:20 Julian Brown
2022-09-28 15:17 ` Tobias Burnus
0 siblings, 1 reply; 4+ messages in thread
From: Julian Brown @ 2022-09-28 13:20 UTC (permalink / raw)
To: gcc-patches; +Cc: Tobias Burnus, Jakub Jelinek, Thomas_Schwinge
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
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;
+}
--
2.29.2
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [PATCH] OpenACC: whole struct vs. component mappings (PR107028)
2022-09-28 13:20 [PATCH] OpenACC: whole struct vs. component mappings (PR107028) Julian Brown
@ 2022-09-28 15:17 ` Tobias Burnus
2022-09-29 12:59 ` Julian Brown
0 siblings, 1 reply; 4+ messages in thread
From: Tobias Burnus @ 2022-09-28 15:17 UTC (permalink / raw)
To: Julian Brown, gcc-patches; +Cc: Jakub Jelinek, Thomas_Schwinge
[-- Attachment #1: Type: text/plain, Size: 3142 bytes --]
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
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [PATCH] OpenACC: whole struct vs. component mappings (PR107028)
2022-09-28 15:17 ` Tobias Burnus
@ 2022-09-29 12:59 ` Julian Brown
2022-09-29 13:09 ` Tobias Burnus
0 siblings, 1 reply; 4+ messages in thread
From: Julian Brown @ 2022-09-29 12:59 UTC (permalink / raw)
To: Tobias Burnus; +Cc: gcc-patches, Jakub Jelinek, Thomas_Schwinge, fortran
[-- Attachment #1: Type: text/plain, Size: 1998 bytes --]
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
[-- Attachment #2: v2-0001-OpenACC-whole-struct-vs.-component-mappings-PR107.patch --]
[-- Type: text/x-patch, Size: 13886 bytes --]
From d0aeea1e93c01d5387c58b8c387018a67e19c5db Mon Sep 17 00:00:00 2001
From: Julian Brown <julian@codesourcery.com>
Date: Tue, 27 Sep 2022 17:39:59 +0000
Subject: [PATCH v2] 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.
---
gcc/gimplify.cc | 176 ++++++++++++++----
.../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 +++
.../gfortran.dg/goacc/mapping-tests-5.f90 | 15 ++
5 files changed, 267 insertions(+), 37 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
create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-5.f90
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 4d032c6bf06..e9fd85b2722 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9861,6 +9861,133 @@ 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 +10365,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 +10496,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 +10537,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 +10733,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;
+}
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
index 00000000000..8df8c5885ad
--- /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
--
2.29.2
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [PATCH] OpenACC: whole struct vs. component mappings (PR107028)
2022-09-29 12:59 ` Julian Brown
@ 2022-09-29 13:09 ` Tobias Burnus
0 siblings, 0 replies; 4+ messages in thread
From: Tobias Burnus @ 2022-09-29 13:09 UTC (permalink / raw)
To: Julian Brown; +Cc: gcc-patches, Jakub Jelinek, Thomas_Schwinge, fortran
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
^ permalink raw reply [flat|nested] 4+ messages in thread
end of thread, other threads:[~2022-09-29 13:10 UTC | newest]
Thread overview: 4+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-09-28 13:20 [PATCH] OpenACC: whole struct vs. component mappings (PR107028) Julian Brown
2022-09-28 15:17 ` Tobias Burnus
2022-09-29 12:59 ` Julian Brown
2022-09-29 13:09 ` Tobias Burnus
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).