From: Tobias Burnus <tobias@codesourcery.com>
To: gcc-patches <gcc-patches@gcc.gnu.org>, Jakub Jelinek <jakub@redhat.com>
Subject: [Patch] OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390)
Date: Mon, 3 Aug 2020 17:37:40 +0200 [thread overview]
Message-ID: <0bbc1a10-2895-ef31-bb9b-95fd0eaac747@codesourcery.com> (raw)
[-- Attachment #1: Type: text/plain, Size: 604 bytes --]
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.)
OK?
Tobias
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
[-- Attachment #2: omp-tmpl.diff --]
[-- Type: text/x-patch, Size: 2488 bytes --]
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<tree> *) 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 <cstdlib>
+#include <type_traits>
+
+template<int Dim> struct V {
+ int version_called;
+
+ template<bool B = (Dim == 0),
+ typename = typename std::enable_if<B>::type>
+ V ()
+ {
+ version_called = 1;
+ }
+
+ template<typename TArg0,
+ typename = typename std::enable_if<(std::is_same<unsigned long,
+ typename std::decay<TArg0>::type>::value)>::type>
+ V (TArg0)
+ {
+ version_called = 2;
+ }
+};
+
+template<int Dim> struct S {
+ V<Dim> 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<long unsigned int>: */
+/* { dg-final { scan-tree-dump-times "__attribute__..omp declare target" 3 "omplower" } } */
next reply other threads:[~2020-08-03 15:37 UTC|newest]
Thread overview: 16+ messages / expand[flat|nested] mbox.gz Atom feed top
2020-08-03 15:37 Tobias Burnus [this message]
2020-08-17 7:17 ` *PING* – " Tobias Burnus
2020-08-25 16:58 ` *PING**2 " Tobias Burnus
2020-08-31 15:53 ` Jakub Jelinek
2020-09-14 13:25 ` Tobias Burnus
2020-09-16 10:36 ` Jakub Jelinek
2020-09-16 23:15 ` Tobias Burnus
2020-09-22 7:11 ` Tobias Burnus
2020-09-22 7:36 ` Jakub Jelinek
2020-09-22 14:11 ` Tobias Burnus
2020-09-22 14:24 ` Jakub Jelinek
2020-09-22 15:39 ` Tobias Burnus
2020-09-23 13:52 ` Tobias Burnus
2020-09-23 14:06 ` Jakub Jelinek
2020-09-23 15:45 ` Tobias Burnus
2020-09-25 11:23 ` Jakub Jelinek
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=0bbc1a10-2895-ef31-bb9b-95fd0eaac747@codesourcery.com \
--to=tobias@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).