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: 1713420 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Received: from sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (P-384) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4NS6KT5dFqz23ys for ; Thu, 8 Dec 2022 06:09:32 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 607AE3833A2B for ; Wed, 7 Dec 2022 19:09:30 +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+incoming=patchwork.ozlabs.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 From patchwork Wed Dec 7 19:13:55 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1713421 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Received: from sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (P-384) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4NS6R86GJqz23pB for ; Thu, 8 Dec 2022 06:14:28 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id D4709392B118 for ; Wed, 7 Dec 2022 19:14:25 +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 22DEF3833A2E; Wed, 7 Dec 2022 19:14:11 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 22DEF3833A2E 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="88732682" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 07 Dec 2022 11:14:09 -0800 IronPort-SDR: hDZKVKYeWhWZaRPIOZWRkuk9BxzZZ87USm1V+oKQBRJX/6sx/beU5hTiBW3n09+wjHY32OPRPS xYZXe45VJWl0bJCQE+ixkYdRKdVLKnZyn2Rz0MaKSPorHLBVwayvRYcWG0JK6Dfq0DrLNBj3vy yUVWk3vkRklJv32/u+qm/gy9UPJPMyTOXjrZGgI8RvHMN7EeDwpR5qd8MjbNsr5RvM3FWrGhvG fpDngHPtQ4Eg9WB1jKoXGy+XwO/IgQwJegJZ4cpALsJmsmh1zkJjtki03Gn17Xu7Nlu9dJS8cN rJc= Date: Wed, 7 Dec 2022 19:13:55 +0000 From: Julian Brown To: Tobias Burnus CC: , Jakub Jelinek , Subject: [PATCH 2/2] OpenMP: Duplicate checking for map clauses in Fortran (PR107214) Message-ID: <20221207191355.2e43ea14@squid.athome> In-Reply-To: <20221207190903.78a6b37f@squid.athome> References: <20221020161414.7430-1-julian@codesourcery.com> <20221207190903.78a6b37f@squid.athome> 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-10.mgc.mentorg.com (147.34.90.210) 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, 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+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" > Hi Julian, > > I had a first quick lock at this patch, I should have a closer look > later. However, I stumbled over the following: > > On 20.10.22 18:14, Julian Brown wrote: > > typedef struct gfc_symbol > > { > > ... > > struct gfc_symbol *old_symbol; > > > > unsigned mark:1, comp_mark:1, data_mark:1, dev_mark:1, > > gen_mark:1; unsigned reduc_mark:1, gfc_new:1; > > > > struct gfc_symbol *tlink; > > > > unsigned equiv_built:1; > > ... > I know that this was the case before, but can you move the mark:1 etc. > after 'tlink'? In that case all bitfields are grouped together. If I > have not miscounted, we have currently 7 bits before and 9 bits after > 'tlink' and grouping them together reduced pointless padding. > > * * * > > + else if (n->sym->mark) > > + gfc_error ("Symbol %qs present on both data and map clauses " > > + "at %L", n->sym->name, &n->where); > > I wonder whether that also rejects the following – which seems to be > valid. The 'map' goes to 'target' and the 'firstprivate' to > 'parallel', cf. OpenMP 5.2, "17.2 Clauses on Combined and Composite > Constructs", [340:3-4 & 12-14]. (BTW: While some fixes went into 5.1 > regarding this section, a likewise wording is already in 5.0.) > > (Testing showed: it give an ICE without the patch and an error with.) ...and this patch avoids the error for combined directives, and reorders the gfc_symbol bitfields. --- This patch adds duplicate checking for OpenMP "map" clauses, taking some cues from the implementation for C in c-typeck.cc:c_finish_omp_clauses (and similar for C++). In addition to the existing use of the "mark" and "comp_mark" bitfields in the gfc_symbol structure, the patch adds several new bits handling duplicate checking within various categories of clause types. If "mark" is being used for map clauses, we need to use different bits for other clauses for cases where "map" and some other clause can refer to the same symbol (e.g. "map(n) shared(n)"). This version of the patch avoids flagging variables that are listed on both map and firstprivate clauses when they are on a combined directive, as they get moved to separate nested directives later (see previous patch in series). Tested with offloading to NVPTX alongside previous patch (and dependencies). OK? 2022-12-06 Julian Brown gcc/fortran/ PR fortran/107214 * gfortran.h (gfc_symbol): Add data_mark, dev_mark, gen_mark and reduc_mark bitfields. * openmp.cc (resolve_omp_clauses): Use above bitfields to improve duplicate clause detection. gcc/testsuite/ PR fortran/107214 * gfortran.dg/gomp/pr107214.f90: New test. From fa6d1e273449aff61833064027fed3787c13121f Mon Sep 17 00:00:00 2001 Message-Id: In-Reply-To: References: From: Julian Brown Date: Tue, 6 Dec 2022 23:10:58 +0000 Subject: [PATCH 2/2] OpenMP: Duplicate checking for map clauses in Fortran (PR107214) This patch adds duplicate checking for OpenMP "map" clauses, taking some cues from the implementation for C in c-typeck.cc:c_finish_omp_clauses (and similar for C++). In addition to the existing use of the "mark" and "comp_mark" bitfields in the gfc_symbol structure, the patch adds several new bits handling duplicate checking within various categories of clause types. If "mark" is being used for map clauses, we need to use different bits for other clauses for cases where "map" and some other clause can refer to the same symbol (e.g. "map(n) shared(n)"). This version of the patch avoids flagging variables that are listed on both map and firstprivate clauses when they are on a combined directive, as they get moved to separate nested directives later (see previous patch in series). Tested with offloading to NVPTX alongside previous patch (and dependencies). OK? 2022-12-06 Julian Brown gcc/fortran/ PR fortran/107214 * gfortran.h (gfc_symbol): Add data_mark, dev_mark, gen_mark and reduc_mark bitfields. * openmp.cc (resolve_omp_clauses): Use above bitfields to improve duplicate clause detection. gcc/testsuite/ PR fortran/107214 * gfortran.dg/gomp/pr107214.f90: New test. --- gcc/fortran/gfortran.h | 32 ++++++--- gcc/fortran/openmp.cc | 73 +++++++++++++++++---- gcc/testsuite/gfortran.dg/gomp/pr107214.f90 | 7 ++ 3 files changed, 90 insertions(+), 22 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/gomp/pr107214.f90 diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index df90ed39bea7..47a7f5552385 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1871,16 +1871,6 @@ typedef struct gfc_symbol gfc_namelist *namelist, *namelist_tail; - /* Change management fields. Symbols that might be modified by the - current statement have the mark member nonzero. Of these symbols, - symbols with old_symbol equal to NULL are symbols created within - the current statement. Otherwise, old_symbol points to a copy of - the old symbol. gfc_new is used in symbol.cc to flag new symbols. - comp_mark is used to indicate variables which have component accesses - in OpenMP/OpenACC directive clauses. */ - struct gfc_symbol *old_symbol; - unsigned mark:1, comp_mark:1, gfc_new:1; - /* The tlink field is used in the front end to carry the module declaration of separate module procedures so that the characteristics can be compared with the corresponding declaration in a submodule. In @@ -1888,6 +1878,28 @@ typedef struct gfc_symbol deferred initialization. */ struct gfc_symbol *tlink; + /* Change management fields. Symbols that might be modified by the + current statement have the mark member nonzero. Of these symbols, + symbols with old_symbol equal to NULL are symbols created within + the current statement. Otherwise, old_symbol points to a copy of + the old symbol. gfc_new is used in symbol.cc to flag new symbols. + comp_mark is used to indicate variables which have component accesses + in OpenMP/OpenACC directive clauses (cf. c-typeck.cc:c_finish_omp_clauses, + map_field_head). + data_mark is used to check duplicate mappings for OpenMP data-sharing + clauses (see firstprivate_head/lastprivate_head in the above function). + dev_mark is used to check duplicate mappings for OpenMP + is_device_ptr/has_device_addr clauses (see is_on_device_head in above + function). + gen_mark is used to check duplicate mappings for OpenMP + use_device_ptr/use_device_addr/private/shared clauses (see generic_head in + above functon). + reduc_mark is used to check duplicate mappings for OpenMP reduction + clauses. */ + struct gfc_symbol *old_symbol; + unsigned mark:1, comp_mark:1, data_mark:1, dev_mark:1, gen_mark:1; + unsigned reduc_mark:1, gfc_new:1; + /* Nonzero if all equivalences associated with this symbol have been processed. */ unsigned equiv_built:1; diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 653c43f79ffb..63a14daa6d7b 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -7135,6 +7135,10 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, continue; n->sym->mark = 0; n->sym->comp_mark = 0; + n->sym->data_mark = 0; + n->sym->dev_mark = 0; + n->sym->gen_mark = 0; + n->sym->reduc_mark = 0; if (n->sym->attr.flavor == FL_VARIABLE || n->sym->attr.proc_pointer || (!code && (!n->sym->attr.dummy || n->sym->ns != ns))) @@ -7203,7 +7207,6 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, && list != OMP_LIST_LASTPRIVATE && list != OMP_LIST_ALIGNED && list != OMP_LIST_DEPEND - && (list != OMP_LIST_MAP || openacc) && list != OMP_LIST_FROM && list != OMP_LIST_TO && (list != OMP_LIST_REDUCTION || !openacc) @@ -7222,10 +7225,43 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next) if (ref->type == REF_COMPONENT) component_ref_p = true; - if ((!component_ref_p && n->sym->comp_mark) - || (component_ref_p && n->sym->mark)) - gfc_error ("Symbol %qs has mixed component and non-component " - "accesses at %L", n->sym->name, &n->where); + if ((list == OMP_LIST_IS_DEVICE_PTR + || list == OMP_LIST_HAS_DEVICE_ADDR) + && !component_ref_p) + { + if (n->sym->gen_mark || n->sym->dev_mark || n->sym->reduc_mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + else + n->sym->dev_mark = 1; + } + else if ((list == OMP_LIST_USE_DEVICE_PTR + || list == OMP_LIST_USE_DEVICE_ADDR + || list == OMP_LIST_PRIVATE + || list == OMP_LIST_SHARED) + && !component_ref_p) + { + if (n->sym->gen_mark || n->sym->dev_mark || n->sym->reduc_mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + else + n->sym->gen_mark = 1; + } + else if (list == OMP_LIST_REDUCTION && !component_ref_p) + { + if (n->sym->gen_mark || n->sym->dev_mark || n->sym->reduc_mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + else + n->sym->reduc_mark = 1; + } + else if ((!component_ref_p && n->sym->comp_mark) + || (component_ref_p && n->sym->mark)) + { + if (openacc) + gfc_error ("Symbol %qs has mixed component and non-component " + "accesses at %L", n->sym->name, &n->where); + } else if (n->sym->mark) gfc_error ("Symbol %qs present on multiple clauses at %L", n->sym->name, &n->where); @@ -7241,31 +7277,44 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, gcc_assert (OMP_LIST_LASTPRIVATE == OMP_LIST_FIRSTPRIVATE + 1); for (list = OMP_LIST_FIRSTPRIVATE; list <= OMP_LIST_LASTPRIVATE; list++) for (n = omp_clauses->lists[list]; n; n = n->next) - if (n->sym->mark) + if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark) { gfc_error ("Symbol %qs present on multiple clauses at %L", n->sym->name, &n->where); - n->sym->mark = 0; + n->sym->data_mark = 0; } + else if (n->sym->mark + && code->op != EXEC_OMP_TARGET_TEAMS + && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE + && code->op != EXEC_OMP_TARGET_TEAMS_LOOP + && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD + && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO + && code->op != EXEC_OMP_TARGET_PARALLEL + && code->op != EXEC_OMP_TARGET_PARALLEL_DO + && code->op != EXEC_OMP_TARGET_PARALLEL_LOOP + && code->op != EXEC_OMP_TARGET_PARALLEL_DO_SIMD + && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD) + gfc_error ("Symbol %qs present on both data and map clauses " + "at %L", n->sym->name, &n->where); for (n = omp_clauses->lists[OMP_LIST_FIRSTPRIVATE]; n; n = n->next) { - if (n->sym->mark) + if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark) gfc_error ("Symbol %qs present on multiple clauses at %L", n->sym->name, &n->where); else - n->sym->mark = 1; + n->sym->data_mark = 1; } for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next) - n->sym->mark = 0; + n->sym->data_mark = 0; for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next) { - if (n->sym->mark) + if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark) gfc_error ("Symbol %qs present on multiple clauses at %L", n->sym->name, &n->where); else - n->sym->mark = 1; + n->sym->data_mark = 1; } for (n = omp_clauses->lists[OMP_LIST_ALIGNED]; n; n = n->next) diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214.f90 new file mode 100644 index 000000000000..25949934e840 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/pr107214.f90 @@ -0,0 +1,7 @@ +! { dg-do compile } + +program p + integer, allocatable :: a + !$omp target map(tofrom: a, a) ! { dg-error "Symbol 'a' present on multiple clauses" } + !$omp end target +end -- 2.29.2