From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 8F5743858D39; Thu, 29 Sep 2022 12:59:34 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 8F5743858D39 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.93,355,1654588800"; d="scan'208,223";a="83728967" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 29 Sep 2022 04:59:31 -0800 IronPort-SDR: /6GxeStuy4d9He9NvxJSZWp8KdAhMLGhzA/U+krbsrHzN8wnkgg/NR0bG5Wy6D8Lu6LqIlapc/ 1hy0ssVAZ1R+N/gppWmWVC7tpvdC2U3JhSuzkDYH4Fr45KgaxX4ZL247w0vHP37n/dRV9vNMFJ GMKPoK4kwvWSxEE+GuE7tQ33Wm1floz61FfuIVUndVcVudO12e30tAe1WGC9c2DNcbXgKc7SyT QU0azHqgtasKV0cvaygPeJ1rMwNPyWrcdWFQ+tLex2iBd9nJjd0hcw4mm7Y6DqMXdu9oZNmanH h5I= Date: Thu, 29 Sep 2022 13:59:11 +0100 From: Julian Brown To: Tobias Burnus CC: , Jakub Jelinek , , Subject: Re: [PATCH] OpenACC: whole struct vs. component mappings (PR107028) Message-ID: <20220929135911.1b2ff3de@squid.athome> In-Reply-To: <66a6261f-cf6f-3a54-b928-eb86ffab4db7@codesourcery.com> References: <20220928132024.64984-1-julian@codesourcery.com> <66a6261f-cf6f-3a54-b928-eb86ffab4db7@codesourcery.com> Organization: Mentor Graphics X-Mailer: Claws Mail 4.0.0 (GTK+ 3.24.31; x86_64-pc-linux-gnu) MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="MP_/vAx67IYOhQ6UJFh0fii46yq" X-ClientProxiedBy: svr-orw-mbx-11.mgc.mentorg.com (147.34.90.211) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-10.4 required=5.0 tests=BAYES_00,GIT_PATCH_0,HEADER_FROM_DIFFERENT_DOMAINS,KAM_DMARC_STATUS,KAM_SHORT,SPF_HELO_PASS,SPF_PASS,TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org List-Id: --MP_/vAx67IYOhQ6UJFh0fii46yq Content-Type: text/plain; charset="UTF-8" Content-Transfer-Encoding: quoted-printable Content-Disposition: inline On Wed, 28 Sep 2022 17:17:30 +0200 Tobias Burnus wrote: > On 28.09.22 15:20, Julian Brown wrote: >=20 > 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? >=20 > OpenACC comments: >=20 > 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 =E2=80=93 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 =E2=80=93 while it compiles *silently* *with* the *patch* applie= d: >=20 > 15 | !$acc enter data copyin(x%A, x%A%i(5), x%A%i(5)) > | ^ > Error: =E2=80=98x.a.i=E2=80=99 appears more than once in map clauses >=20 > 15 | !$acc enter data copyin(x%A, x%A%i(5), x%A%i(4)) > | ^ > Error: =E2=80=98x.a.i=E2=80=99 appears more than once in map clauses >=20 > 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 --MP_/vAx67IYOhQ6UJFh0fii46yq Content-Type: text/x-patch Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="v2-0001-OpenACC-whole-struct-vs.-component-mappings-PR107.patch" >From d0aeea1e93c01d5387c58b8c387018a67e19c5db Mon Sep 17 00:00:00 2001 From: Julian Brown 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 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 *groups, + hash_map *grpmap) +{ + int i; + omp_mapping_group *grp; + hash_set *seen_components = NULL; + hash_set *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 (); + if (!shown_error) + shown_error = new hash_set (); + 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 *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 + +#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 + +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 --MP_/vAx67IYOhQ6UJFh0fii46yq--