From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1534) id B98A239960FD; Thu, 17 Sep 2020 17:03:46 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org B98A239960FD DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1600362226; bh=OSUOBaafG3xSv0ARVUp14RU66xyt5Iu/fuRCHo45RJ0=; h=From:To:Subject:Date:From; b=UOSTY11mDdn/KvSAlFZSdQWnVAcHE4MPTzMKmPy6U5rDkcVOZ8HwetGM2e8w4nd26 16T25nxUNh2/XmNN6x2hbfpkUV7l1c9Qt6/XIdzUFiRiJR/DqHewmo69pReRbWi1Nq bjUlVHQhm5Oky8of6ZWS475cuwzza2QS96cObhi0= Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Tobias Burnus To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-10] OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390) X-Act-Checkin: gcc X-Git-Author: Tobias Burnus X-Git-Refname: refs/heads/devel/omp/gcc-10 X-Git-Oldrev: 37788b29ddfb7686ccb71debad347f68eebd6293 X-Git-Newrev: ef509d1985aa53a8c0875c25ad4050ea807be10e Message-Id: <20200917170346.B98A239960FD@sourceware.org> Date: Thu, 17 Sep 2020 17:03:46 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Thu, 17 Sep 2020 17:03:47 -0000 https://gcc.gnu.org/g:ef509d1985aa53a8c0875c25ad4050ea807be10e commit ef509d1985aa53a8c0875c25ad4050ea807be10e Author: Tobias Burnus Date: Thu Sep 17 18:45:00 2020 +0200 OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390) Submitted to mainline at https://gcc.gnu.org/pipermail/gcc-patches/2020-September/554142.html 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. Diff: --- gcc/ChangeLog.omp | 6 +++ gcc/omp-offload.c | 50 ++++++++++++++++++++---- libgomp/ChangeLog.omp | 6 +++ libgomp/testsuite/libgomp.c++/pr96390.C | 49 +++++++++++++++++++++++ libgomp/testsuite/libgomp.c-c++-common/pr96390.c | 26 ++++++++++++ 5 files changed, 129 insertions(+), 8 deletions(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 63ccd8a3a79..f666abe13ac 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,9 @@ +2020-09-17 Tobias Burnus + + PR middle-end/96390 + * omp-offload.c (omp_discover_declare_target_tgt_fn_r): Handle + alias nodes. + 2020-09-17 Tobias Burnus Backport from mainline diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index 9c0ca643a2f..72be733a38d 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -198,21 +198,55 @@ 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 (); + 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/ChangeLog.omp b/libgomp/ChangeLog.omp index 82e7cc2d941..297c6055293 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,9 @@ +2020-09-17 Tobias Burnus + + PR middle-end/96390 + * testsuite/libgomp.c++/pr96390.C: New test. + * testsuite/libgomp.c-c++-common/pr96390.c: New test. + 2020-09-17 Tobias Burnus Backport from mainline 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" } } */ 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..3857bf3348d --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c @@ -0,0 +1,26 @@ +/* { dg-run } */ +/* { dg-additional-options "-O0 -fdump-tree-omplower" } */ + +#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" } } */