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 59B5F385782F for ; Tue, 22 Sep 2020 07:11:56 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 59B5F385782F Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=mentor.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Tobias_Burnus@mentor.com IronPort-SDR: w2LrbxJI0pxAsBsmazVnAl3xEfg7jIuWVme71z7UkR5g350jF86ka+uqK7N56To43Q9bI97yjU KI651STN4eubd9EDaUr0ez2t+wFl6jkjKZ80CV96NEAauIEDAlpQKqTC5zg5p8b9Su3aF/qPdn zpXUPzB9OnWzzaRSeWPuDijj1JOZI6/2hzONMAxC7pRJlol1VKDBQKQx3ULy7SuNrpIEa0NhUA KN6ZYk2BYwFtKSL/rOC0U4mCWdqUmZS0QGqyMbQiVzuwgO6e482EzaFhfVO554vQMnf3NkH/gm y3M= X-IronPort-AV: E=Sophos;i="5.77,289,1596528000"; d="diff'?scan'208";a="53163223" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 21 Sep 2020 23:11:54 -0800 IronPort-SDR: bqTD9qonLrXBd8YprUaORTGy0zGwHfKWcAw5Mf4/6IcUGoK1IyqpjxDjZdzi+a9qMFK4sTKGKq ZAie+HhE8Kp+pxteNDlEIQM7SL9NY6Rlt5gs0H8riolclpYNAbIP6npduJA6TXceC2gR1JA3fO lz6+Z+9tjUI4RHF28Hd0GAuInASJDWeQZbtnQlI0gaNfPIEtvH0xKt5K9sPnWBQREUs9m/iY2C ysNK2iXWUypF0UQiVqSvW6oPUIgvJUNJBwDH2e3ZaTGWjC0bdjzwltgNRilB2WravQK6t2sfOI vis= Subject: Re: [Patch] OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390) To: Jakub Jelinek , Tobias Burnus CC: Jan Hubicka , gcc-patches References: <0bbc1a10-2895-ef31-bb9b-95fd0eaac747@codesourcery.com> <20200831155311.GS18149@tucnak> <20200916103647.GV21814@tucnak> From: Tobias Burnus Message-ID: Date: Tue, 22 Sep 2020 09:11:48 +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: Content-Type: multipart/mixed; boundary="------------4012E6F22D19C773F8AF50FF" Content-Language: en-US X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-08.mgc.mentorg.com (139.181.222.8) 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, 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: Tue, 22 Sep 2020 07:11:58 -0000 --------------4012E6F22D19C773F8AF50FF Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: quoted-printable Hi Honza & Jakub, @Honza: please look at the decl alias handling of the actual patch. a minor update =E2=80=93 testing only on gcn did turn out to be insufficien= t: nvptx does not support alias, cf. PR 97102 + 97106; hence xfailed. I also had a typo in one 'dg-do run' but as -O0 is set explicitly, 'dg-do run' does not make sense. (no dg-do =3D run once, dg-do run =3D run with multiple options, dg-do compile =3D compile only). Hence, I have now removed the dg-do line. On 9/17/20 1:15 AM, Tobias Burnus wrote: > Hi Honza =E2=80=93 some input would be really helpful! > > Hi Jakub =E2=80=93 updated version below. > > On 9/16/20 12:36 PM, Jakub Jelinek wrote: >> I think you want Honza on this primarily, I'm always lost >> in the cgraph alias code. > (Likewise as this thread shows) >>> + while (node->alias_target) >>> + node =3D symtab_node::get (node->alias_target); >>> + node =3D node->ultimate_alias_target (); >> I think the above is either you walk the aliases yourself, or use >> ultimate_alias_target, but not both. > > I think we need to distinguish between: > * aliases which end up with the same symbol name > and are stored in the ref_list; example: cpp_implicit_alias. > * aliases like with the alias attribute, which is handled > via alias_target and have different names. > > Just experimentally: > * The 'while (node->alias_target)' properly resolves the > attribute testcase (libgomp.c-c++-common/pr96390.c). > Here, ultimate_alias_target () does not help as > node->analyzed =3D=3D 0. > > * The 'node->ultimate_alias_target ()' works for the > cpp_implicit_alias case (libgomp.c++/pr96390.C). > Just looking at the alias target does not help as in this > case, alias_target =3D=3D NULL. > >> And the second thing is, I'm not sure how the aliases behave if the >> ultimate alias target is properly marked as omp declare target, but some >> of the aliases are not. The offloaded code will still call the alias, >> so do we somehow arrange for the aliases to be also emitted into the >> offloading LTO IL? >> [...] I wonder if the aliases that are needed shouldn't >> be marked node->offloadable and have "omp declare target" attribute >> added >> for them too. > > Done now. > > Okay =E2=80=93 or do we find more issues? > > 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 --------------4012E6F22D19C773F8AF50FF Content-Type: text/x-patch; charset="UTF-8"; name="omp-tmpl-v6.diff" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="omp-tmpl-v6.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 | 51 ++++++++++++++++++++---- libgomp/testsuite/libgomp.c++/pr96390.C | 49 +++++++++++++++++++++++ libgomp/testsuite/libgomp.c-c++-common/pr96390.c | 26 ++++++++++++ 3 files changed, 118 insertions(+), 8 deletions(-) diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index 32c2485abd4..222ebff1d1e 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -196,21 +196,56 @@ 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 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)); + tree decl = *tp; symtab_node *node = symtab_node::get (*tp); if (node != NULL) { - node->offloadable = 1; + /* First, find final FUNCTION_DECL; find final alias target and there + ensure alias like cpp_implicit_alias are resolved by calling + ultimate_alias_target; the latter does not resolve alias_target as + node->analyzed = 0. */ + symtab_node *orig_node = node; + while (node->alias_target) + node = symtab_node::get (node->alias_target); + node = node->ultimate_alias_target (); + decl = node->decl; + + if (omp_declare_target_fn_p (decl) + || lookup_attribute ("omp declare target host", + DECL_ATTRIBUTES (decl))) + return NULL_TREE; + if (ENABLE_OFFLOADING) g->have_offload = true; + + /* Now mark original node and all alias targets for offloading. */ + node->offloadable = 1; + if (orig_node != node) + { + tree id = get_identifier ("omp declare target"); + while (orig_node->alias_target) + { + orig_node = orig_node->ultimate_alias_target (); + orig_node->offloadable = 1; + DECL_ATTRIBUTES (orig_node->decl) + = tree_cons (id, NULL_TREE, + DECL_ATTRIBUTES (orig_node->decl)); + orig_node = symtab_node::get (orig_node->alias_target); + } + } } + else if (omp_declare_target_fn_p (decl) + || lookup_attribute ("omp declare target host", + DECL_ATTRIBUTES (decl))) + return NULL_TREE; + + tree id = get_identifier ("omp declare target"); + 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" } } */ --------------4012E6F22D19C773F8AF50FF--