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 D40CC385800F; Wed, 10 Mar 2021 14:20:51 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org D40CC385800F 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: w5J+XJFAhxTsrmCe0siEL9qRZME3cw2+ublILGuVXWhf3J331wA1/2dl7N4KiSdje5tawOqsgT YV4qMPc8HjebjdIU63IuUam9yHrLmTwMAD5Za7HxonEJG3TpTxt6HMlI9abiqtj39XyaEa8KiB MgWz56osiDZXlAmhhmCV/NJCVseGanbCml97+JiDJY2N3S6dygpilZkedT3D2r+3tc14ehd84N hk7jGmvCU6dlaKqbfc/qV7CWrrw1x7fQU8EJmQqqtITWCLTMwL3THegwlohtcw9tGIqSrF9wgV gqM= X-IronPort-AV: E=Sophos;i="5.81,237,1610438400"; d="diff'?scan'208";a="59075607" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 10 Mar 2021 06:20:50 -0800 IronPort-SDR: W7mMiUxBoAiuHDUtSN75ePk+bETCHGQ2ZKKL3ShJpPCKIshQhHb7hIr4KnJd6FgtlJ5p/Y8gOF JG9vmaiTR0whCV8mc6yFsiY6GzqLwtGbUrWsWeLIUJFosoqvmZkWzMO//LLSkuEytA7qigf8IH f5yYOuLEUVLIJ4qryvLyU49Wm20pUsv1xS0vsmKmlR+pQ22BGMYgTglDZGasfsbEvcSOi/M+wE pP/VPyBTLHE9phXWE+rBKSCaVCG/CgsZbnYIa2v6Dacdz2xnsOGI5eaMxBIvJF4QfBoSmMt6Qo Z6s= To: gcc-patches , fortran , Jakub Jelinek From: Tobias Burnus Subject: [Patch] OpenMP: Fix 'omp declare target' handling for vars [PR99509] Message-ID: <29dc3864-050b-a83d-4b64-3982e233bd4f@codesourcery.com> Date: Wed, 10 Mar 2021 15:20:42 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:78.0) Gecko/20100101 Thunderbird/78.8.0 MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="------------027ECC1F43CC04107D94EBA0" 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, 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: Wed, 10 Mar 2021 14:20:54 -0000 --------------027ECC1F43CC04107D94EBA0 Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: quoted-printable The C/C++ FE sets for an 'omp declare target' ... 'omp end declare target' 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 target l= ink' attribute. Unfortunately, adding 'omp declare target' comes too late as the varpool has been generated. I am not sure what's the best way, but the following seems to work. 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 --------------027ECC1F43CC04107D94EBA0 Content-Type: text/x-patch; charset="UTF-8"; name="omp-decl-tgt.diff" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="omp-decl-tgt.diff" OpenMP: Fix 'omp declare target' handling for vars [PR99509] C/C++ use 'omp declare target implicit' when they call varpool_node::get_create; make sure that a later 'omp declare target' (but not 'omp declare target link') is honored. gcc/ChangeLog: PR c++/99509 * varpool.c (varpool_node::get_create): Re-check whether 'omp declare target'. libgomp/ChangeLog: PR c++/99509 * testsuite/libgomp.c-c++-common/declare_target-1.c: New test. gcc/varpool.c | 18 +++++++++++++++++- .../libgomp.c-c++-common/declare_target-1.c | 22 ++++++++++++++++++++++ 2 files changed, 39 insertions(+), 1 deletion(-) diff --git a/gcc/varpool.c b/gcc/varpool.c index 4830df5..5ae6778 100644 --- a/gcc/varpool.c +++ b/gcc/varpool.c @@ -145,7 +145,23 @@ varpool_node::get_create (tree decl) varpool_node *node = varpool_node::get (decl); gcc_checking_assert (VAR_P (decl)); if (node) - return node; + { + /* This happens with C/C++ when 'omp declare target implicit' is set + which is only later replaced by 'omp declate target' or '... link'. */ + if (!node->offloadable + && (flag_openacc || flag_openmp) + && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))) + { + node->offloadable = 1; + if (ENABLE_OFFLOADING && !DECL_EXTERNAL (decl)) + { + g->have_offload = true; + if (!in_lto_p) + vec_safe_push (offload_vars, decl); + } + } + return node; + } node = varpool_node::create_empty (); node->decl = decl; 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 0000000..c5670df --- /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; +} --------------027ECC1F43CC04107D94EBA0--