Hello
OpenMP 5.0 has a new feature for implicitly marking variables and functions that
are referenced in the initializers of static variables and functions that are
already marked 'declare target'. Support was added in the commit 'openmp:
Implement discovery of implicit declare target to clauses'
(dc703151d4f4560e647649506d5b4ceb0ee11e90). However, this does not work with
non-constant C++ initializers, where the initializers can contain references to
other (non-constant) variables and function calls.
The C++ front-end stores the initialization information in the static_aggregates
list (with the variable decl in the TREE_VALUE of an entry and the
initialization in TREE_PURPOSE) rather than in TREE_INITIAL(var_decl). I have
added an extra function in omp-offload.cpp to walk the variable initialiser
trees in static_aggregates, and added a call to it from the FE shortly before
the initializations are emitted. I have also added a testcase to ensure that the
implicitly marked variables/functions can be referenced in offloaded code.
The libgomp tests have been run with offloading to a Nvidia card with no
regressions, and I have also bootstrapped the compiler with no offloading on
x86-64. Okay for trunk?
Thanks
Kwok
commit d2c8c5bd2826851b727e93a8ea2141596e50a621
Author: Kwok Cheung Yeung <k...@codesourcery.com>
Date: Wed Oct 28 07:13:14 2020 -0700
openmp: Implicitly add 'declare target' directives for dynamic static
initializers in C++
2020-10-28 Kwok Cheung Yeung <k...@codesourcery.com>
cp/
* decl2.c: Include omp-offload.h
(c_parse_final_cleanups): Call omp_mark_target_static_initializers.
* omp-offload.c (omp_discover_declare_target_var_r): Add all static
variables to worklist.
(omp_discover_implicit_declare_target): Check that worklist items
that are variable declarations have an initialization expression
before walking.
(omp_mark_target_static_initializers): New.
* omp-offload.h (omp_mark_target_static_initializers): New prototype.
libgomp/
* testsuite/libgomp.c++/declare_target-3.C: New.
diff --git a/gcc/cp/decl2.c b/gcc/cp/decl2.c
index 2f0d637..b207d58 100644
--- a/gcc/cp/decl2.c
+++ b/gcc/cp/decl2.c
@@ -48,6 +48,7 @@ along with GCC; see the file COPYING3. If not see
#include "intl.h"
#include "c-family/c-ada-spec.h"
#include "asan.h"
+#include "omp-offload.h"
/* Id for dumping the raw trees. */
int raw_dump_id;
@@ -4970,6 +4971,12 @@ c_parse_final_cleanups (void)
/* Make sure the back end knows about all the variables. */
write_out_vars (vars);
+ /* Mark functions and variables in static initializers as
+ 'omp declare target' if the initialized variable is marked
+ as such. */
+ if (flag_openmp)
+ omp_mark_target_static_initializers (vars);
+
/* Set the line and file, so that it is obviously not from
the source file. */
input_location = locus_at_end_of_parsing;
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 3e9c31d..8ecc181 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -296,7 +296,7 @@ omp_discover_declare_target_var_r (tree *tp, int
*walk_subtrees, void *data)
DECL_ATTRIBUTES (*tp)
= remove_attribute ("omp declare target link", DECL_ATTRIBUTES
(*tp));
}
- if (TREE_STATIC (*tp) && DECL_INITIAL (*tp))
+ if (TREE_STATIC (*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);
@@ -348,7 +348,7 @@ omp_discover_implicit_declare_target (void)
while (!worklist.is_empty ())
{
tree decl = worklist.pop ();
- if (VAR_P (decl))
+ if (VAR_P (decl) && DECL_INITIAL (decl))
walk_tree_without_duplicates (&DECL_INITIAL (decl),
omp_discover_declare_target_var_r,
&worklist);
@@ -363,6 +363,33 @@ omp_discover_implicit_declare_target (void)
}
}
+void
+omp_mark_target_static_initializers (tree vars)
+{
+ tree node;
+ auto_vec<tree> worklist;
+
+ for (node = vars; node; node = TREE_CHAIN (node))
+ if (omp_declare_target_var_p (TREE_VALUE (node)))
+ worklist.safe_push (TREE_VALUE (node));
+
+ while (!worklist.is_empty ())
+ {
+ tree decl = worklist.pop ();
+
+ if (!VAR_P (decl) || !TREE_STATIC (decl))
+ continue;
+
+ for (node = vars; node; node = TREE_CHAIN (node))
+ if (TREE_VALUE (node) == decl)
+ {
+ walk_tree_without_duplicates (&TREE_PURPOSE (node),
+ omp_discover_declare_target_var_r,
+ &worklist);
+ break;
+ }
+ }
+}
/* Create new symbols containing (address, size) pairs for global variables,
marked with "omp declare target" attribute, as well as addresses for the
diff --git a/gcc/omp-offload.h b/gcc/omp-offload.h
index 0809189..a23a4d3 100644
--- a/gcc/omp-offload.h
+++ b/gcc/omp-offload.h
@@ -31,5 +31,6 @@ extern GTY(()) vec<tree, va_gc> *offload_vars;
extern void omp_finish_file (void);
extern void omp_discover_implicit_declare_target (void);
+extern void omp_mark_target_static_initializers (tree vars);
#endif /* GCC_OMP_DEVICE_H */
diff --git a/libgomp/testsuite/libgomp.c++/declare_target-3.C
b/libgomp/testsuite/libgomp.c++/declare_target-3.C
new file mode 100644
index 0000000..c545613
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/declare_target-3.C
@@ -0,0 +1,36 @@
+/* { dg-do run } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#include <stdlib.h>
+
+int f()
+{
+ return 54321;
+}
+
+int g()
+{
+ return 3333;
+}
+
+static int x = f() + 7777;
+static int y = g() + x + 12345;
+
+#pragma omp declare target(y)
+
+int main()
+{
+ int err = 0;
+ #pragma omp target map(from:err)
+ {
+ err |= x != 62098;
+ err |= y != 77776;
+ err |= f() != 54321;
+ err |= g() != 3333;
+ }
+ if (err)
+ abort ();
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump-times "__attribute__\\(\\(omp declare
target\\)\\)" 2 "gimple" } } */
