From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 51EB53858001; Mon, 15 Mar 2021 08:17:59 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 51EB53858001 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Tobias_Burnus@mentor.com IronPort-SDR: /2ba97nPBLF+2U63A4/KIKA+9kyt8Fq7kPXkpDo1JgmV8HLGBJ34klybnYjAlXrRuSiGTZWR03 N8td4LHhIeg6j5pvX5XHmAhjtdzngq6g8meJ0Y6OJEoRacvkDhebZj4FWjBkw15sEbntL5eAXh ks4IbFIxd2Yw4aOAWOOx+lXE11WB8mI3RySKxTZgPCCOelw1o8VIhmSLVYvQ1tyVJxsDh5a+qh gHEtvrcZf0dtKaaNr2k3wRetnKwLRwb7cFPkZ4+vqQMdllxIVFiFEPjfYjg/v8Xm0lhzd6XKin UvQ= X-IronPort-AV: E=Sophos;i="5.81,249,1610438400"; d="diff'?scan'208";a="59073436" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 15 Mar 2021 00:17:57 -0800 IronPort-SDR: wAqrauivQ6QyR3+xzAWcSdPj66vMJ3wqWDJC7ZpbK/AubH6ulKEO3fdQrk3YCJ3QUinbJZORbc oEl9MhEAhB3YTsMMdZ/WSd6wCaglWcFjMiddznpD99cVeuEZEkObMBwI5SKEPKO0GhbyuP3Lxf M+Cj11nNBowu48aDf1q393LPbdO8qVwSSJDKTuNexmD//ozmCqUaHniClFXqXPAXJJR32lWstD TxeDXVXy14FHB457eQJpuvl3FbSjJfO754nL/4ETRBduiKxEgCglgGJQb7KSrPnYVSVy5Vbqhv 1js= Subject: Re: [Patch] OpenMP: Fix 'omp declare target' handling for vars [PR99509] To: Jakub Jelinek CC: gcc-patches , fortran References: <29dc3864-050b-a83d-4b64-3982e233bd4f@codesourcery.com> <20210312122049.GT745611@tucnak> From: Tobias Burnus Message-ID: <93bf9c6b-1c60-017c-422c-5b4112f88509@codesourcery.com> Date: Mon, 15 Mar 2021 09:17:47 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:78.0) Gecko/20100101 Thunderbird/78.8.1 MIME-Version: 1.0 In-Reply-To: <20210312122049.GT745611@tucnak> Content-Type: multipart/mixed; boundary="------------D644167EED6BFF09A38ABCE6" Content-Language: en-US X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.1 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, NICE_REPLY_A, RCVD_IN_DNSWL_NONE, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) on server2.sourceware.org X-BeenThere: fortran@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Fortran mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Mon, 15 Mar 2021 08:18:01 -0000 --------------D644167EED6BFF09A38ABCE6 Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: quoted-printable On 12.03.21 13:20, Jakub Jelinek wrote: > On Wed, Mar 10, 2021 at 03:20:42PM +0100, Tobias Burnus wrote: >> The C/C++ FE sets for an 'omp declare target' ... 'omp end declare targe= t' >> the attribute 'omp declare target implicit'. >> >> That's later processed (for C++) in decl.c - which remove that attribute >> and either keeps and explicit 'omp declare target' or 'omp declare targe= t link' >> attribute. >> >> Unfortunately, adding 'omp declare target' comes too late as the varpool >> has been generated. >> >> Looking attributes on every get rather than just once is IMHO too expens= ive. >> For explicit declare target, we use: >> ... >> (e.g. from cp/parser.c). >> So, I think it would be better to do the same thing I concur. I was already wondering whether it should be in the FE or not, thinking about some separation between FE and ME. But given the precedent and the performance issue, it surely makes sense to solve it in c*/*decl.c. Done now. OK for mainline? Tobias ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 M=C3=BCnchen R= egistergericht M=C3=BCnchen HRB 106955, Gesch=C3=A4ftsf=C3=BChrer: Thomas H= eurung, Frank Th=C3=BCrauf --------------D644167EED6BFF09A38ABCE6 Content-Type: text/x-patch; charset="UTF-8"; name="omp-decl-tgt-v2.diff" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="omp-decl-tgt-v2.diff" OpenMP: Fix 'omp declare target' handling for vars [PR99509] For variables with 'declare target' attribute, varpool_node::get_create marks variables as offload; however, if the node already exists, it is not updated. C/C++ may tag decl with 'declare target implicit', which may only be after varpool creation turned into 'declare target' or 'declare target link'; in this case, the tagging has to happen in the FE. gcc/c/ChangeLog: PR c++/99509 * c-decl.c (finish_decl): For 'omp declare target implicit' vars, ensure that the varpool node is marked as offloadable. gcc/cp/ChangeLog: PR c++/99509 * decl.c (cp_finish_decl): For 'omp declare target implicit' vars, ensure that the varpool node is marked as offloadable. libgomp/ChangeLog: PR c++/99509 * testsuite/libgomp.c-c++-common/declare_target-1.c: New test. gcc/c/c-decl.c | 22 +++++++++++++++++++--- gcc/cp/decl.c | 21 ++++++++++++++++++--- .../libgomp.c-c++-common/declare_target-1.c | 22 ++++++++++++++++++++++ 3 files changed, 59 insertions(+), 6 deletions(-) diff --git a/gcc/c/c-decl.c b/gcc/c/c-decl.c index b559ed5d76a..3b2241bfd97 100644 --- a/gcc/c/c-decl.c +++ b/gcc/c/c-decl.c @@ -58,6 +58,9 @@ along with GCC; see the file COPYING3. If not see #include "c-family/name-hint.h" #include "c-family/known-headers.h" #include "c-family/c-spellcheck.h" +#include "context.h" /* For 'g'. */ +#include "omp-general.h" +#include "omp-offload.h" /* For offload_vars. */ #include "tree-pretty-print.h" @@ -5658,9 +5661,22 @@ finish_decl (tree decl, location_t init_loc, tree init, DECL_ATTRIBUTES (decl)) && !lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (decl))) - DECL_ATTRIBUTES (decl) - = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, DECL_ATTRIBUTES (decl)); + { + DECL_ATTRIBUTES (decl) + = tree_cons (get_identifier ("omp declare target"), + NULL_TREE, DECL_ATTRIBUTES (decl)); + symtab_node *node = symtab_node::get (decl); + if (node != NULL) + { + node->offloadable = 1; + if (ENABLE_OFFLOADING) + { + g->have_offload = true; + if (is_a (node)) + vec_safe_push (offload_vars, decl); + } + } + } } /* This is the last point we can lower alignment so give the target the diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c index 9c7f6e59bbb..1b671cec9a3 100644 --- a/gcc/cp/decl.c +++ b/gcc/cp/decl.c @@ -53,7 +53,9 @@ along with GCC; see the file COPYING3. If not see #include "asan.h" #include "gcc-rich-location.h" #include "langhooks.h" +#include "context.h" /* For 'g'. */ #include "omp-general.h" +#include "omp-offload.h" /* For offload_vars. */ /* Possible cases of bad specifiers type used by bad_specifiers. */ enum bad_spec_place { @@ -8176,9 +8178,22 @@ cp_finish_decl (tree decl, tree init, bool init_const_expr_p, DECL_ATTRIBUTES (decl)) && !lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (decl))) - DECL_ATTRIBUTES (decl) - = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, DECL_ATTRIBUTES (decl)); + { + DECL_ATTRIBUTES (decl) + = tree_cons (get_identifier ("omp declare target"), + NULL_TREE, DECL_ATTRIBUTES (decl)); + symtab_node *node = symtab_node::get (decl); + if (node != NULL) + { + node->offloadable = 1; + if (ENABLE_OFFLOADING) + { + g->have_offload = true; + if (is_a (node)) + vec_safe_push (offload_vars, decl); + } + } + } } /* This is the last point we can lower alignment so give the target the diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c b/libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c new file mode 100644 index 00000000000..c5670dfb7db --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c @@ -0,0 +1,22 @@ +/* PR c++/99509 */ + +#pragma omp declare target +int data[] = {5}; +#pragma omp end declare target + +static inline int +foo (int idx) +{ + return data[idx]; +} + +int +main () +{ + int i = -1; + #pragma omp target map(from:i) + i = foo(0); + if (i != 5) + __builtin_abort (); + return 0; +} --------------D644167EED6BFF09A38ABCE6--