From patchwork Wed Dec 7 19:09:03 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 61669 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id E5EF3392B108 for ; Wed, 7 Dec 2022 19:09:34 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id 68F6638493EA; Wed, 7 Dec 2022 19:09:15 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 68F6638493EA 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,225,1665475200"; d="scan'208,223";a="92016974" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 07 Dec 2022 11:09:13 -0800 IronPort-SDR: zLt3L7EV9DEsOsXyLSimENt6DeorHmPYdm2uji6c33f1ZH04XhIojXzkaBGwSZoXDYGdRnIwH4 BKX1fKCqBsPsaMNYHL2fgOvE7jZSX/cML4gUbb4RWrA0XrYIElcxQIxfXgwmK2127DSrFOEYIV s80KK13tL2veNxOSJuDkdTmJ1ZK8QQD4014GktBP+nFr10d2aAGsgR9G7nyaSdZyHzRYQ0R8he K0N6figJpgxjmMrn929TIZtdDqx/fJW3GmJxbfmxclFVXqkqD+IEBTH3FDOm4jouqEgiRSfNWF xIA= Date: Wed, 7 Dec 2022 19:09:03 +0000 From: Julian Brown To: Tobias Burnus CC: , Jakub Jelinek , Subject: [PATCH 1/2] OpenMP/Fortran: Combined directives with map/firstprivate of same symbol Message-ID: <20221207190903.78a6b37f@squid.athome> In-Reply-To: References: <20221020161414.7430-1-julian@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 X-ClientProxiedBy: SVR-ORW-MBX-07.mgc.mentorg.com (147.34.90.207) 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, 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: , Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" On Wed, 26 Oct 2022 12:39:39 +0200 Tobias Burnus wrote: > The ICE seems to be because gcc/fortran/trans-openmp.cc's > gfc_split_omp_clauses mishandles this as the dump shows the following: > > #pragma omp target firstprivate(a) map(tofrom:a) > #pragma omp parallel firstprivate(a) > > * * * > > In contrast, for the C testcase: > > void foo(int x) { > #pragma omp target parallel for simd map(x) firstprivate(x) > for (int k = 0; k < 1; ++k) > x = 1; > } > > the dump is as follows, which seems to be sensible: > > #pragma omp target map(tofrom:x) > #pragma omp parallel firstprivate(x) > #pragma omp for nowait > #pragma omp simd First, here's a patch to address this bit... This patch fixes a case where a combined directive (e.g. "!$omp target parallel ...") contains both a map and a firstprivate clause for the same variable. When the combined directive is split into two nested directives, the outer "target" gets the "map" clause, and the inner "parallel" gets the "firstprivate" clause, like so: !$omp target parallel map(x) firstprivate(x) --> !$omp target map(x) !$omp parallel firstprivate(x) ... When there is no map of the same variable, the firstprivate is distributed to both directives, e.g. for 'y' in: !$omp target parallel map(x) firstprivate(y) --> !$omp target map(x) firstprivate(y) !$omp parallel firstprivate(y) ... This is not a recent regression, but appears to fix a long-standing ICE. (The included testcase is based on one by Tobias.) Tested with offloading to NVPTX, alongside previously-posted patches (in review or approved but waiting for other patches), i.e.: OpenMP/OpenACC: Rework clause expansion and nested struct handling OpenMP/OpenACC: Refine condition for when map clause expansion happens OpenMP: Pointers and member mappings and the patch following. OK? 2022-12-06 Julian Brown gcc/fortran/ * trans-openmp.cc (gfc_add_firstprivate_if_unmapped): New function. (gfc_split_omp_clauses): Call above. libgomp/ * testsuite/libgomp.fortran/combined-directive-splitting-1.f90: New test. From c66db363066913ae4939f2aa706427338b109d71 Mon Sep 17 00:00:00 2001 Message-Id: From: Julian Brown Date: Tue, 6 Dec 2022 12:18:33 +0000 Subject: [PATCH 1/2] OpenMP/Fortran: Combined directives with map/firstprivate of same symbol This patch fixes a case where a combined directive (e.g. "!$omp target parallel ...") contains both a map and a firstprivate clause for the same variable. When the combined directive is split into two nested directives, the outer "target" gets the "map" clause, and the inner "parallel" gets the "firstprivate" clause, like so: !$omp target parallel map(x) firstprivate(x) --> !$omp target map(x) !$omp parallel firstprivate(x) ... When there is no map of the same variable, the firstprivate is distributed to both directives, e.g. for 'y' in: !$omp target parallel map(x) firstprivate(y) --> !$omp target map(x) firstprivate(y) !$omp parallel firstprivate(y) ... This is not a recent regression, but appears to fix a long-standing ICE. (The included testcase is based on one by Tobias.) Tested with offloading to NVPTX, alongside previously-posted patches (in review or approved but waiting for other patches), i.e.: OpenMP/OpenACC: Rework clause expansion and nested struct handling OpenMP/OpenACC: Refine condition for when map clause expansion happens OpenMP: Pointers and member mappings and the patch following. OK? 2022-12-06 Julian Brown gcc/fortran/ * trans-openmp.cc (gfc_add_firstprivate_if_unmapped): New function. (gfc_split_omp_clauses): Call above. libgomp/ * testsuite/libgomp.fortran/combined-directive-splitting-1.f90: New test. --- gcc/fortran/trans-openmp.cc | 37 ++++++++++++++++- .../combined-directive-splitting-1.f90 | 41 +++++++++++++++++++ 2 files changed, 76 insertions(+), 2 deletions(-) create mode 100644 libgomp/testsuite/libgomp.fortran/combined-directive-splitting-1.f90 diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index e39f7b1cb273..c61cd1bf55de 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -6121,6 +6121,39 @@ gfc_add_clause_implicitly (gfc_omp_clauses *clauses_out, } } +/* Kind of opposite to above, add firstprivate to CLAUSES_OUT if it is mapped + in CLAUSES_IN's FIRSTPRIVATE list but not its MAP list. */ + +static void +gfc_add_firstprivate_if_unmapped (gfc_omp_clauses *clauses_out, + gfc_omp_clauses *clauses_in) +{ + gfc_omp_namelist *n = clauses_in->lists[OMP_LIST_FIRSTPRIVATE]; + gfc_omp_namelist **tail = NULL; + + for (; n != NULL; n = n->next) + { + gfc_omp_namelist *n2 = clauses_out->lists[OMP_LIST_MAP]; + for (; n2 != NULL; n2 = n2->next) + if (n->sym == n2->sym) + break; + if (n2 == NULL) + { + gfc_omp_namelist *dup = gfc_get_omp_namelist (); + *dup = *n; + dup->next = NULL; + if (!tail) + { + tail = &clauses_out->lists[OMP_LIST_FIRSTPRIVATE]; + while (*tail && (*tail)->next) + tail = &(*tail)->next; + } + *tail = dup; + tail = &(*tail)->next; + } + } +} + static void gfc_free_split_omp_clauses (gfc_code *code, gfc_omp_clauses *clausesa) { @@ -6504,8 +6537,8 @@ gfc_split_omp_clauses (gfc_code *code, simd and masked/master. Put it on the outermost of those and duplicate on parallel and teams. */ if (mask & GFC_OMP_MASK_TARGET) - clausesa[GFC_OMP_SPLIT_TARGET].lists[OMP_LIST_FIRSTPRIVATE] - = code->ext.omp_clauses->lists[OMP_LIST_FIRSTPRIVATE]; + gfc_add_firstprivate_if_unmapped (&clausesa[GFC_OMP_SPLIT_TARGET], + code->ext.omp_clauses); if (mask & GFC_OMP_MASK_TEAMS) clausesa[GFC_OMP_SPLIT_TEAMS].lists[OMP_LIST_FIRSTPRIVATE] = code->ext.omp_clauses->lists[OMP_LIST_FIRSTPRIVATE]; diff --git a/libgomp/testsuite/libgomp.fortran/combined-directive-splitting-1.f90 b/libgomp/testsuite/libgomp.fortran/combined-directive-splitting-1.f90 new file mode 100644 index 000000000000..e662a2bd3b20 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/combined-directive-splitting-1.f90 @@ -0,0 +1,41 @@ +module m + integer :: a = 1 + !$omp declare target enter(a) +end module m + +module m2 +contains +subroutine bar() + use m + implicit none + !$omp declare target + a = a + 5 +end subroutine bar +end module m2 + +program p + use m + use m2 + implicit none + integer :: b, i + + !$omp target parallel do map(always, tofrom: a) firstprivate(a) + do i = 1, 1 + a = 7 + call bar() + if (a /= 7) error stop 1 + a = a + 8 + end do + if (a /= 6) error stop 2 + + b = 3 + !$omp target parallel do map(always, tofrom: a) firstprivate(b) + do i = 1, 1 + a = 3 + call bar () + if (a /= 8) error stop 3 + a = a + b + end do + if (a /= 11) error stop 4 +end program p + -- 2.29.2