From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from us-smtp-1.mimecast.com (us-smtp-2.mimecast.com [205.139.110.61]) by sourceware.org (Postfix) with ESMTP id 943C2384C005 for ; Mon, 31 Aug 2020 15:53:57 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 943C2384C005 Received: from mimecast-mx01.redhat.com (mimecast-mx01.redhat.com [209.132.183.4]) (Using TLS) by relay.mimecast.com with ESMTP id us-mta-326-YRXF-deKOzeKOWI8XsLtMA-1; Mon, 31 Aug 2020 11:53:53 -0400 X-MC-Unique: YRXF-deKOzeKOWI8XsLtMA-1 Received: from smtp.corp.redhat.com (int-mx03.intmail.prod.int.phx2.redhat.com [10.5.11.13]) (using TLSv1.2 with cipher AECDH-AES256-SHA (256/256 bits)) (No client certificate requested) by mimecast-mx01.redhat.com (Postfix) with ESMTPS id 799409392FE; Mon, 31 Aug 2020 15:53:16 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-113-115.ams2.redhat.com [10.36.113.115]) by smtp.corp.redhat.com (Postfix) with ESMTPS id 155B37EB77; Mon, 31 Aug 2020 15:53:15 +0000 (UTC) Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.15.2/8.15.2) with ESMTP id 07VFrDYU015414; Mon, 31 Aug 2020 17:53:13 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.15.2/8.15.2/Submit) id 07VFrBe7015413; Mon, 31 Aug 2020 17:53:11 +0200 Date: Mon, 31 Aug 2020 17:53:11 +0200 From: Jakub Jelinek To: Tobias Burnus , Jan Hubicka Cc: gcc-patches Subject: Re: [Patch] OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390) Message-ID: <20200831155311.GS18149@tucnak> Reply-To: Jakub Jelinek References: <0bbc1a10-2895-ef31-bb9b-95fd0eaac747@codesourcery.com> MIME-Version: 1.0 In-Reply-To: <0bbc1a10-2895-ef31-bb9b-95fd0eaac747@codesourcery.com> User-Agent: Mutt/1.11.3 (2019-02-01) X-Scanned-By: MIMEDefang 2.79 on 10.5.11.13 X-Mimecast-Spam-Score: 0.002 X-Mimecast-Originator: redhat.com Content-Type: text/plain; charset=utf-8 Content-Transfer-Encoding: 8bit Content-Disposition: inline X-Spam-Status: No, score=-11.9 required=5.0 tests=BAYES_00, DKIMWL_WL_HIGH, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, RCVD_IN_MSPIKE_H3, RCVD_IN_MSPIKE_WL, SPF_HELO_NONE, 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: Mon, 31 Aug 2020 15:53:59 -0000 On Mon, Aug 03, 2020 at 05:37:40PM +0200, Tobias Burnus wrote: > It turned out that the omp_discover_declare_target_tgt_fn_r > discovered all nodes – but as it tagged the C++ alias nodes > and not the streamed-out nodes, no device function was created > and one got link errors if offloading devices were configured. > (Only with -O0 as otherwise inlining happened.) > > (Testcase is based on a sollve_vv testcase which in turn was > based on an LLVM bugreport.) > 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 > cpp_implicit_alias nodes. > > libgomp/ChangeLog: > > PR middle-end/96390 > * testsuite/libgomp.c++/pr96390.C: New test. > > gcc/omp-offload.c | 8 ++++++ > libgomp/testsuite/libgomp.c++/pr96390.C | 49 +++++++++++++++++++++++++++++++++ > 2 files changed, 57 insertions(+) > > diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c > index 32c2485abd4..4aef7dbea6c 100644 > --- a/gcc/omp-offload.c > +++ b/gcc/omp-offload.c > @@ -207,6 +207,14 @@ omp_discover_declare_target_tgt_fn_r (tree *tp, int *walk_subtrees, void *data) > symtab_node *node = symtab_node::get (*tp); > if (node != NULL) > { > + if (node->cpp_implicit_alias) > + { > + node = node->get_alias_target (); > + if (!omp_declare_target_fn_p (node->decl)) > + ((vec *) data)->safe_push (node->decl); > + DECL_ATTRIBUTES (node->decl) > + = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (node->decl)); > + } > node->offloadable = 1; > if (ENABLE_OFFLOADING) > g->have_offload = true; Sorry for the review delay. I don't see what is special on cpp_implicit_alias here, compared to any other aliases. So, I wonder if the code shouldn't do: tree decl = *tp; symtab_node *node = symtab_node::get (decl); if (node != NULL) { symtab_node *anode = node->ultimate_alias_target (); if (anode && anode != node) { decl = anode->decl; gcc_assert (TREE_CODE (decl) == FUNCTION_DECL); if (omp_declare_target_fn_p (*tp) || lookup_attribute ("omp declare target host", DECL_ATTRIBUTES (decl))) return NULL_TREE; node = anode; } } 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)); if (node != NULL) { node->offloadable = 1; if (ENABLE_OFFLOADING) g->have_offload = true; } Otherwise, if we have say: void foo () { } void bar () __attribute__((alias ("foo"))); void baz () __attribute__((alias ("bar"))); int main () { #pragma omp target baz (); } we won't mark foo as being declare target. Though, maybe that is not enough and we need to mark all the aliases from node to the ultimate alias that way (so perhaps instead iterate one get_alias_target (if node->alias) by one and mark all the decls the way the code marks right now (i.e. pushes those non-DECL_EXTERNAL with DECL_SAVED_TREE for further processing and add attributes to all of them? Jakub