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 6C1DF3857C46 for ; Wed, 23 Sep 2020 15:45:19 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 6C1DF3857C46 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: U1/IJaSb+A6gz2r66OEu0QBsoilLHht1Lnu+kvvbcXpq8FQ2uY3ArKCKYANwxOFlI2C698jSzb 9w0guQeaiQZyQwfVY/US9lOGmc0uHGp2yM2xADXv6NHs91FbWUsIOR7ffO/ws66R0u1lcfaAvp 8Au/3NzDglsjC/+0eaaLHvAk/JWXotIfVghtpXx1LH8n6K8wPwMvM4kRT7SEJ0ilWSkd4oMqwI RUZ+PZuhUFevwwDD5kiMCRa3P4joESOELmVjCWXzMfCdF+Xx194Sexg+JKYZ6aiUAhMqe2qB6S 9ag= X-IronPort-AV: E=Sophos;i="5.77,293,1596528000"; d="diff'?scan'208";a="53232815" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 23 Sep 2020 07:45:18 -0800 IronPort-SDR: HdAvUIW4TPPMS+CMqSdqA8xZZvNui9mmjv2vweAf3qGqK1VIjstGsFUNxOuexJacoSZ2yFsR65 Fsi1mewFRWPi+jRQR1+VWZs5qmGKBine2LPeZJBtEMCndhqxPWqvKsBTScJNFs2QFaR8K/YQtH 5durhd3KPchq+oJWZQqMiFxffILSkFmco/3Ug+yMzw/XYMo51R+AipH+vzuSk2XKFGRO8UvjCv WWIW47cYMMNt491f1TSOUcW3Wwf3oasPuN4KhtCguFC4lxyAIHDXI6eGsyVLnDeaGb/3ToHixE OU4= Subject: Re: [Patch] OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390) To: Jakub Jelinek , Tobias Burnus CC: Jan Hubicka , gcc-patches References: <20200831155311.GS18149@tucnak> <20200916103647.GV21814@tucnak> <20200922073659.GV2176@tucnak> <07508a46-2ea3-ac17-a398-a0f4dd56c556@codesourcery.com> <20200922142433.GZ2176@tucnak> <54423ac9-a651-246f-3570-feb976f4b67e@codesourcery.com> <20200923140642.GH2176@tucnak> From: Tobias Burnus Message-ID: <165d3742-8097-7dc6-621e-efefcf476d3f@codesourcery.com> Date: Wed, 23 Sep 2020 17:45:12 +0200 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:68.0) Gecko/20100101 Thunderbird/68.12.0 MIME-Version: 1.0 In-Reply-To: <20200923140642.GH2176@tucnak> Content-Type: multipart/mixed; boundary="------------4B71E0E48FEF1C0C931249CE" Content-Language: en-US X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, NICE_REPLY_A, 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: 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: , X-List-Received-Date: Wed, 23 Sep 2020 15:45:21 -0000 --------------4B71E0E48FEF1C0C931249CE Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: quoted-printable On 9/23/20 4:06 PM, Jakub Jelinek wrote: > What I really meant was: I did now something based on this. >> + gcc_assert (node->alias && node->analyzed); I believe from previous testing that node->analyzed is 0 for the testcase at hand =E2=80=94 and, hence, ultimate_alias_target() did not walk node->target_alias for the testcase at hand. (=E2=86=92 symtab_node::ultimate_alias_target_1) Hence, I didn't include it in the assert. But I can re-check. Tobias ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstra=C3=9Fe 201, 80634 M=C3=BCnch= en / Germany Registergericht M=C3=BCnchen HRB 106955, Gesch=C3=A4ftsf=C3=BChrer: Thomas = Heurung, Alexander Walter --------------4B71E0E48FEF1C0C931249CE Content-Type: text/x-patch; charset="UTF-8"; name="omp-tmpl-v9.diff" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="omp-tmpl-v9.diff" OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390) gcc/ChangeLog: PR middle-end/96390 * omp-offload.c (omp_discover_declare_target_tgt_fn_r): Handle alias nodes. libgomp/ChangeLog: PR middle-end/96390 * testsuite/libgomp.c++/pr96390.C: New test. * testsuite/libgomp.c-c++-common/pr96390.c: New test. gcc/omp-offload.c | 44 ++++++++++++++++++--- libgomp/testsuite/libgomp.c++/pr96390.C | 49 ++++++++++++++++++++++++ libgomp/testsuite/libgomp.c-c++-common/pr96390.c | 26 +++++++++++++ 3 files changed, 113 insertions(+), 6 deletions(-) diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index 32c2485abd4..dd5e9a2c9b2 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -196,21 +196,53 @@ omp_declare_target_var_p (tree decl) static tree omp_discover_declare_target_tgt_fn_r (tree *tp, int *walk_subtrees, void *data) { - if (TREE_CODE (*tp) == FUNCTION_DECL - && !omp_declare_target_fn_p (*tp) - && !lookup_attribute ("omp declare target host", DECL_ATTRIBUTES (*tp))) + if (TREE_CODE (*tp) == FUNCTION_DECL) { + tree decl = *tp; tree id = get_identifier ("omp declare target"); - if (!DECL_EXTERNAL (*tp) && DECL_SAVED_TREE (*tp)) - ((vec *) data)->safe_push (*tp); - DECL_ATTRIBUTES (*tp) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (*tp)); symtab_node *node = symtab_node::get (*tp); if (node != NULL) { + while (node->alias_target) + { + if (!omp_declare_target_fn_p (node->decl) + && !lookup_attribute ("omp declare target host", + DECL_ATTRIBUTES (node->decl))) + { + node->offloadable = 1; + DECL_ATTRIBUTES (node->decl) + = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (node->decl)); + } + node = symtab_node::get (node->alias_target); + } + symtab_node *new_node = node->ultimate_alias_target (); + decl = new_node->decl; + while (node != new_node) + { + if (!omp_declare_target_fn_p (node->decl) + && !lookup_attribute ("omp declare target host", + DECL_ATTRIBUTES (node->decl))) + { + node->offloadable = 1; + DECL_ATTRIBUTES (node->decl) + = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (node->decl)); + } + gcc_assert (node->alias); + node = node->get_alias_target (); + } node->offloadable = 1; if (ENABLE_OFFLOADING) g->have_offload = true; } + if (omp_declare_target_fn_p (decl) + || lookup_attribute ("omp declare target host", + DECL_ATTRIBUTES (decl))) + return NULL_TREE; + + if (!DECL_EXTERNAL (decl) && DECL_SAVED_TREE (decl)) + ((vec *) data)->safe_push (decl); + DECL_ATTRIBUTES (decl) = tree_cons (id, NULL_TREE, + DECL_ATTRIBUTES (decl)); } else if (TYPE_P (*tp)) *walk_subtrees = 0; diff --git a/libgomp/testsuite/libgomp.c++/pr96390.C b/libgomp/testsuite/libgomp.c++/pr96390.C new file mode 100644 index 00000000000..8c770ecb80c --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/pr96390.C @@ -0,0 +1,49 @@ +/* { dg-additional-options "-O0 -fdump-tree-omplower" } */ +/* { dg-xfail-if "PR 97106/PR 97102 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */ + +#include +#include + +template struct V { + int version_called; + + template::type> + V () + { + version_called = 1; + } + + template::type>::value)>::type> + V (TArg0) + { + version_called = 2; + } +}; + +template struct S { + V v; +}; + +int +main () +{ + int version_set[2] = {-1, -1}; + +#pragma omp target map(from: version_set[0:2]) + { + S<0> s; + version_set[0] = s.v.version_called; + V<1> v2((unsigned long) 1); + version_set[1] = v2.version_called; + } + + if (version_set[0] != 1 || version_set[1] != 2) + abort (); + return 0; +} + +/* "3" for S<0>::S, V<0>::V<>, and V<1>::V: */ +/* { dg-final { scan-tree-dump-times "__attribute__..omp declare target" 3 "omplower" } } */ diff --git a/libgomp/testsuite/libgomp.c-c++-common/pr96390.c b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c new file mode 100644 index 00000000000..692bd730069 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c @@ -0,0 +1,26 @@ +/* { dg-additional-options "-O0 -fdump-tree-omplower" } */ +/* { dg-xfail-if "PR 97102/PR 97106 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */ + +#ifdef __cplusplus +extern "C" { +#endif + +int foo () { return 42; } +int bar () __attribute__((alias ("foo"))); +int baz () __attribute__((alias ("bar"))); + +#ifdef __cplusplus +} +#endif + + +int +main () +{ + int n; + #pragma omp target map(from:n) + n = baz (); + if (n != 42) + __builtin_abort (); +} +/* { dg-final { scan-tree-dump-times "__attribute__..omp declare target" 1 "omplower" } } */ --------------4B71E0E48FEF1C0C931249CE--