On 9/22/20 9:36 AM, Jakub Jelinek wrote:
On Tue, Sep 22, 2020 at 09:11:48AM +0200, Tobias Burnus wrote:
Okay – or do we find more issues?
I'm afraid so.
We will slowly converge, hopefully ;-)
Consider:
int v;
#pragma omp declare target to (v)
void foo (void) { v++; }
void bar (void) __attribute__((alias ("foo")));
#pragma omp declare target to (bar)
void baz (void) __attribute__((alias ("foo")));
void qux (void) {
etc. – I did not convert this into a testcase. Should I?
Do you spot something more? Or is it now fine? (It passes on gcn +
xfailed on nvptx.)
Tobias
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander
Walter
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 | 38 +++++++++++++++---
libgomp/testsuite/libgomp.c++/pr96390.C | 49 ++++++++++++++++++++++++
libgomp/testsuite/libgomp.c-c++-common/pr96390.c | 26 +++++++++++++
3 files changed, 107 insertions(+), 6 deletions(-)
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 32c2485abd4..a1b12bbbb8f 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -196,21 +196,47 @@ 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 decl = *tp;
tree id = get_identifier ("omp declare target");
- if (!DECL_EXTERNAL (*tp) && DECL_SAVED_TREE (*tp))
- ((vec<tree> *) data)->safe_push (*tp);
- DECL_ATTRIBUTES (*tp) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (*tp));
symtab_node *node = symtab_node::get (*tp);
if (node != NULL)
{
+ while (node->alias_target)
+ {
+ node = node->ultimate_alias_target ();
+ if (!omp_declare_target_fn_p (node->decl)
+ && !lookup_attribute ("omp declare target host",
+ DECL_ATTRIBUTES (node->decl)))
+ {
+ node->offloadable = 1;
+ DECL_ATTRIBUTES (node->decl)
+ = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (node->decl));
+ }
+ 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;
+
node->offloadable = 1;
if (ENABLE_OFFLOADING)
g->have_offload = true;
}
+ else if (omp_declare_target_fn_p (decl)
+ || lookup_attribute ("omp declare target host",
+ DECL_ATTRIBUTES (decl)))
+ return NULL_TREE;
+
+ if (!DECL_EXTERNAL (decl) && DECL_SAVED_TREE (decl))
+ ((vec<tree> *) 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 <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" } } */
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" } } */
