From patchwork Wed Sep 28 13:20:24 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1508 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:adf:f2ce:0:0:0:0:0 with SMTP id d14csp165742wrp; Wed, 28 Sep 2022 06:22:11 -0700 (PDT) X-Google-Smtp-Source: AMsMyM73qxYXhrt5zbWEkoGyb+2bhuH6NPUw+ju/Jfxg+z65mdU9zjFnBYPHCz2BZOM4Ru+Kh21l X-Received: by 2002:a05:6402:1e96:b0:451:129e:1a2c with SMTP id f22-20020a0564021e9600b00451129e1a2cmr33673190edf.60.1664371331035; Wed, 28 Sep 2022 06:22:11 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1664371331; cv=none; d=google.com; s=arc-20160816; b=DBBTQZv1Lt3fJ8gNOS2LhibG53TKaPtkFlFvQraN6mMD7ytbMYhEYpgJ2Am7aEbm8x 92vFOfYRx5Ta9QCqhqhUNBuxwb+5HF7fDHKeyOEQFkctu9+yOBK3FlFUuoXoxUGRf6iR Zo4C4+VlAVDfPrNYS8zfJ4nuv+fzCwXCU+upUMe6cgUDztoDldWuvJ1CXWpxHge260aY mlYJpvH8PUz7npO2Qk6YIWoTocHdYCJ4QkI7EcIZyKUa5hkCfqIv2m0LjVrTbU3ccCZj V+JJ2FbiRX45LGKUlQG5igVvSw9VmO9+sE8xhwcuujKc18hHMwkETzrRk/nwtTKSI38y qeiA== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:cc:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:content-transfer-encoding :mime-version:message-id:date:subject:to:from:ironport-sdr :dmarc-filter:delivered-to; bh=tQ1S5yAWSJHbvpsS+ghttEp8x0XB3+X9XJtkaHH0x4g=; b=CPCSKCv4wyXnNcFz/Wb7h259iqYKfhj2GWvmdtjI3JjH+8xt2JRAlbe7H0H5BNfvsk EObpSBWYRWbepK/Yx+rkPNLGnrW5hVVBCkPbB4yZOoHPu4T4k4b6Z2+5PhFrRt74bAFE o2tE8HsQvS0if3Gi1WXyLC2mHrq1rza0Ce7eFIImzgaec4IN8ccfjRrUAVbDUGYvwSts 3/BqXCm/8WcmY9inP0NUIubdXI2QgQ/SWhVGuJc/Y1rirya2XQXOTBIUVR6DhxM85VJg M2SHDKtC4OOT7cd0daZnlxehlzktjKN3hmXtuO4/+p3cDX7tk59ulkA/SJ0y/+/CTUtH 38mQ== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from sourceware.org (ip-8-43-85-97.sourceware.org. [8.43.85.97]) by mx.google.com with ESMTPS id b6-20020aa7c6c6000000b00454412dc7absi1802815eds.259.2022.09.28.06.22.10 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 28 Sep 2022 06:22:11 -0700 (PDT) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) client-ip=8.43.85.97; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 019663858283 for ; Wed, 28 Sep 2022 13:22:09 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 1F9FA385802A for ; Wed, 28 Sep 2022 13:20:48 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 1F9FA385802A 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,352,1654588800"; d="scan'208";a="83654469" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 28 Sep 2022 05:20:39 -0800 IronPort-SDR: T9JsL/Hjkuq4d+eeWXHh9MLuGkDXv5GuZx6bGAjLb1yBRoaeZhaiyjyFDwbXq59r0V68rDw305 m8hgeHjS5/61t7br2OcpBU9pVJ3ogcgb7EGmFjrHarV+4NfYzrfyxDwA50NNtT0T+RwqumIzFl zd/BZNOZMGyYQdCyezJT+O6EzVlBOjX5ykE5HyOyBcZA+FvXmOhiKBfLk2DRe/jSRiGvFEyGLU scqPAhQ/4IyxjGQ7KM83zVJmk783Fc/Ur0jrIdl9sWVvgt43FwFbUVDO5L4wL6pQGp6yU8ykyl ywE= From: Julian Brown To: Subject: [PATCH] OpenACC: whole struct vs. component mappings (PR107028) Date: Wed, 28 Sep 2022 06:20:24 -0700 Message-ID: <20220928132024.64984-1-julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-14.mgc.mentorg.com (139.181.222.14) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-11.8 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 X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Cc: Jakub Jelinek , Thomas_Schwinge@mentor.com, Tobias Burnus Errors-To: gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: =?utf-8?q?INBOX?= X-GMAIL-THRID: =?utf-8?q?1745219832449219112?= X-GMAIL-MSGID: =?utf-8?q?1745219832449219112?= 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 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 *groups, + hash_map *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 *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; +}