From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 89F3C3857C71; Mon, 12 Dec 2022 15:19:56 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 89F3C3857C71 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.96,238,1665475200"; d="diff'?scan'208";a="89371868" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 12 Dec 2022 07:19:53 -0800 IronPort-SDR: 7A2ENlRhP50LFktQ//2kQfDlZkxnG4xSDvkBUGPokIQk8u5G5HVoLujYND+VxtroeAVn3nrRRS f1Eu3plXoD1DxVR8LoS1i6bE39EFDX25ZYi/iwSSkrw/PO55+EL1F2mX4zH3valEuoi3icPO++ YXdkVvTXzC/cJQ7XR96Dlnilv1437Mb49qv0blPktYm9dXjuHhPaqn34qXkSkzcv8O6CVxQFHS 3e9Ezv0b3njP4dJVcPpDsex0yPMkDMa7PVYAwmC6itME0SLZ5Zeuhnx89mY1FzLRcvDadwiHzo Ick= Date: Mon, 12 Dec 2022 15:19:40 +0000 From: Julian Brown To: Tobias Burnus CC: , , Jakub Jelinek , Chung-Lin Tang Subject: Re: [PATCH v5 2/4] OpenMP/OpenACC: Rework clause expansion and nested struct handling Message-ID: <20221212151940.55c2abaf@squid.athome> In-Reply-To: <81d1d3c5-4f37-21ca-f6e5-bd75bc197c2b@codesourcery.com> References: <8fcf3df1b40ea77cbb8088962cbcdf6935d2ded3.1666088224.git.julian@codesourcery.com> <65824aa6-be83-3c74-871e-6571008f2d25@codesourcery.com> <20221207151657.7340469c@squid.athome> <81d1d3c5-4f37-21ca-f6e5-bd75bc197c2b@codesourcery.com> Organization: Siemens Embedded X-Mailer: Claws Mail 4.1.1git7 (GTK 3.24.34; x86_64-pc-linux-gnu) MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="MP_/fEHC/asrqtGpl.GkwtQRhKE" X-ClientProxiedBy: svr-orw-mbx-08.mgc.mentorg.com (147.34.90.208) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-10.5 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_/fEHC/asrqtGpl.GkwtQRhKE Content-Type: text/plain; charset="US-ASCII" Content-Transfer-Encoding: 7bit Content-Disposition: inline On Wed, 7 Dec 2022 17:13:42 +0100 Tobias Burnus wrote: > I think we need to distinguish: > > #pragma omp target enter data map(to: s.w[:10]) > > from > > #pragma omp target map(tofrom: s.arr[:20]) > s.arr[0] = 5; > > As in the latter case 's' gets implicitly mapped and then applies to > the base pointer 's.arr' of 's.arr[:20]'. While in the former case, > only the pointee gets mapped without the pointer 's.arr' (and, hence, > there is also no pointer attachment). Here's an incremental patch that fixes the mapping behaviour in that case. This is to be applied on top of the approved (but not committed) parent patch: https://gcc.gnu.org/pipermail/gcc-patches/2022-October/603792.html and also the unreviewed patch posted here (ping?): https://gcc.gnu.org/pipermail/gcc-patches/2022-November/607543.html though it might actually make more sense to commit the three patches squashed together. Tested with offloading to NVPTX. OK? Thanks, Julian --MP_/fEHC/asrqtGpl.GkwtQRhKE Content-Type: text/x-patch Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="implicitly-map-base-pointer-1.diff" commit abb1e04f9ef93221ecd4292f43cc1ea901843766 Author: Julian Brown Date: Thu Dec 8 13:31:01 2022 +0000 OpenMP: implicitly map base pointer for array-section pointer components Following from discussion in: https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570075.html and: https://gcc.gnu.org/pipermail/gcc-patches/2022-December/608100.html and also upstream OpenMP issue 342, this patch changes mapping for array sections of pointer components on compute regions like this: #pragma omp target map(s.ptr[0:10]) { ...use of 's'... } so the base pointer 's.ptr' is implicitly mapped, and thus pointer attachment happens. This is subtly different in the "enter data" case, e.g: #pragma omp target enter data map(s.ptr[0:10]) if 's.ptr' (or the whole of 's') is not present on the target before the directive is executed, the array section is copied to the target but pointer attachment does *not* take place, since 's' (or 's.ptr') is not mapped implicitly for "enter data". To get a pointer attachment with "enter data", you can do, e.g: #pragma omp target enter data map(s.ptr, s.ptr[0:10]) #pragma omp target { ...implicit use of 's'... } That is, once the attachment has happened, implicit mapping of 's' and uses of 's.ptr[...]' work correctly in the target region. ChangeLog 2022-12-12 Julian Brown gcc/ * gimplify.cc (omp_accumulate_sibling_list): Don't require explicitly-mapped base pointer for compute regions. gcc/testsuite/ * c-c++-comon/gomp/target-implicit-map-2.c: Update expected scan output. libgomp/ * testsuite/libgomp.c-c++-common/target-implicit-map-2.c: Fix missing "free". * testsuite/libgomp.c-c++-common/target-implicit-map-3.c: New test. * testsuite/libgomp.c-c++-common/target-map-zlas-1.c: New test. * testsuite/libgomp.c/target-22.c: Remove explicit base pointer mappings. diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 34cac30d7d92..a8dd298559e8 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -10617,6 +10617,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type, poly_int64 cbitpos; tree ocd = OMP_CLAUSE_DECL (grp_end); bool openmp = !(region_type & ORT_ACC); + bool target = (region_type & ORT_TARGET) != 0; tree *continue_at = NULL; while (TREE_CODE (ocd) == ARRAY_REF) @@ -10721,9 +10722,9 @@ omp_accumulate_sibling_list (enum omp_region_type region_type, } /* For OpenMP semantics, we don't want to implicitly allocate - space for the pointer here. A FRAGILE_P node is only being - created so that omp-low.cc is able to rewrite the struct - properly. + space for the pointer here for non-compute regions (e.g. "enter + data"). A FRAGILE_P node is only being created so that + omp-low.cc is able to rewrite the struct properly. For references (to pointers), we want to actually allocate the space for the reference itself in the sorted list following the struct node. @@ -10731,6 +10732,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type, mapping of the attachment point, but not otherwise. */ if (*fragile_p || (openmp + && !target && attach_detach && TREE_CODE (TREE_TYPE (ocd)) == POINTER_TYPE && !OMP_CLAUSE_ATTACHMENT_MAPPING_ERASED (grp_end))) @@ -11043,6 +11045,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type, if (*fragile_p || (openmp + && !target && attach_detach && TREE_CODE (TREE_TYPE (ocd)) == POINTER_TYPE && !OMP_CLAUSE_ATTACHMENT_MAPPING_ERASED (grp_end))) diff --git a/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c index 5ba1d7efe08d..222272df5b1e 100644 --- a/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c +++ b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c @@ -49,4 +49,5 @@ main (void) /* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:a \[len: [0-9]+\]\[implicit\]\)} "gimple" } } */ -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(struct:a \[len: 1\]\) map\(alloc:a\.ptr \[len: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(attach:a\.ptr \[bias: 0\]\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(struct:a \[len: 1\]\) map\(alloc:a\.ptr \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(attach:a\.ptr \[bias: 0\]\)} "gimple" } } */ +/* { dg-final { scan-tree-dump-not {map\(struct:a \[len: 1\]\) map\(alloc:a\.ptr \[len: 0\]\)} "gimple" } } */ diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c index 974a9786c3f6..4c49cd091c38 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c +++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c @@ -42,5 +42,7 @@ main (void) #pragma omp target exit data map(from:a.ptr, a.ptr[:N]) + free (a.ptr); + return 0; } diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c new file mode 100644 index 000000000000..81a7752685c5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c @@ -0,0 +1,50 @@ +#include + +#define N 10 + +struct S +{ + int a, b; + int *ptr; + int c, d; +}; + +int +main (void) +{ + struct S a; + a.ptr = (int *) malloc (sizeof (int) * N); + + for (int i = 0; i < N; i++) + a.ptr[i] = 0; + + #pragma omp target enter data map(to: a.ptr) + #pragma omp target enter data map(to: a.ptr[:N]) + + #pragma omp target + for (int i = 0; i < N; i++) + a.ptr[i] += 1; + + #pragma omp target update from(a.ptr[:N]) + + for (int i = 0; i < N; i++) + if (a.ptr[i] != 1) + abort (); + + #pragma omp target map(a.ptr[:N]) + for (int i = 0; i < N; i++) + a.ptr[i] += 1; + + #pragma omp target update from(a.ptr[:N]) + + for (int i = 0; i < N; i++) + if (a.ptr[i] != 2) + abort (); + + #pragma omp target exit data map(release: a.ptr[:N]) + #pragma omp target exit data map(release: a.ptr) + + free (a.ptr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-zlas-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-zlas-1.c new file mode 100644 index 000000000000..1ec0c9a0d5f9 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-zlas-1.c @@ -0,0 +1,36 @@ +#include + +#define N 10 + +struct S +{ + int a, b; + int *ptr; + int c, d; +}; + +int +main (void) +{ + struct S a; + a.ptr = (int *) malloc (sizeof (int) * N); + + for (int i = 0; i < N; i++) + a.ptr[i] = 0; + + #pragma omp target enter data map(to: a.ptr[:N]) + + #pragma omp target map(a, a.ptr[:0]) + for (int i = 0; i < N; i++) + a.ptr[i] += 1; + + #pragma omp target exit data map(from: a.ptr[:N]) + + for (int i = 0; i < N; i++) + if (a.ptr[i] != 1) + abort (); + + free (a.ptr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-22.c b/libgomp/testsuite/libgomp.c/target-22.c index 492744ad0efd..aad8a0a09df7 100644 --- a/libgomp/testsuite/libgomp.c/target-22.c +++ b/libgomp/testsuite/libgomp.c/target-22.c @@ -21,8 +21,7 @@ main () s.v.b = a + 16; s.w = c + 3; int err = 0; - #pragma omp target map (to: s.w, s.v.b, s.u, s.s) \ - map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \ + #pragma omp target map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \ map (tofrom:s.s[3:3]) \ map (from: s.w[z:4], err) private (i) { --MP_/fEHC/asrqtGpl.GkwtQRhKE--