From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id BB5963857C53 for ; Mon, 3 Aug 2020 15:37:46 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org BB5963857C53 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: O2Xgy0skpjKf2IaELiv3skYWfYg9mSIqhK41PoxiErRNsnMthDU/rcPjAQ9pwgcVjx2fW/WfhC yt+fl3DPUKm04tvk4HPQx+wLhDtsMu5aIo22vJvyLM5A/LXESZ+9SxjbpNnhsvYWiVoo9NhLzg weyJslp5h/hJRARGxGBLkVXSq4YZn8/BvVUfAuQ8dkItU7gZKLnAG7sGIlLPLsxQeJ68oZW+ws CeAUaZIgUPwdCgi25B4fbDXvx4shFg5SbqXO1weUz19QG0Occ2RpyBoOynX1o71go7jDe9RD8D 6Qo= X-IronPort-AV: E=Sophos;i="5.75,430,1589270400"; d="diff'?scan'208";a="53720436" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa1.mentor.iphmx.com with ESMTP; 03 Aug 2020 07:37:45 -0800 IronPort-SDR: FPdHyhsKWdSDt6b7iRzLI5QYM1SFFdgkeHD0TDjhyBEmRkYihnyY8jSQ3zN7Wjh9CBALB+LTW/ SBJIvY73qQiWJ51jD8js/ANNIyIAN5J9egKsVsyO51WNO/sif+56SIfn9av0Rzvbz70d9PO4rQ VhwihUEFmtjGelDMvodmPsOzvpu7IE3ksRjiW/1RR4dAyhKhnbgEPECBOFvQSyCLp+3xDcAHcC TmhDndmdMXzU8Lz9SjlTk4LY/yyRWglBhHOMCGKFP6lP47CkMdggG8tE7/pKZHu8T6eE8KkFdM y4U= To: gcc-patches , Jakub Jelinek From: Tobias Burnus Subject: [Patch] OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390) Message-ID: <0bbc1a10-2895-ef31-bb9b-95fd0eaac747@codesourcery.com> Date: Mon, 3 Aug 2020 17:37:40 +0200 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:68.0) Gecko/20100101 Thunderbird/68.10.0 MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="------------0EF7541295D8DA749A1B3859" 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=-12.9 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: 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, 03 Aug 2020 15:37:48 -0000 --------------0EF7541295D8DA749A1B3859 Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: quoted-printable It turned out that the omp_discover_declare_target_tgt_fn_r discovered all nodes =E2=80=93 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.) OK? 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 --------------0EF7541295D8DA749A1B3859 Content-Type: text/x-patch; charset="UTF-8"; name="omp-tmpl.diff" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="omp-tmpl.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 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; diff --git a/libgomp/testsuite/libgomp.c++/pr96390.C b/libgomp/testsuite/libgomp.c++/pr96390.C new file mode 100644 index 00000000000..098cb103919 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/pr96390.C @@ -0,0 +1,49 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O0 -fdump-tree-omplower" } */ + +#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" } } */ --------------0EF7541295D8DA749A1B3859--