This patch adds support for running constructors and destructors for
static (file-scope) aggregates for C++ objects which are marked with
"declare target" directives on OpenMP offload targets.
At present, space is allocated on the target for such aggregates, but
nothing ever constructs them properly, so they end up zero-initialised.
Tested with offloading to AMD GCN. I will apply to the og12 branch
shortly.
ChangeLog
2023-03-27 Julian Brown <jul...@codesourcery.com>
gcc/cp/
* decl2.cc (priority_info): Add omp_tgt_initializations_p and
omp_tgt_destructions_p.
(start_objects, start_static_storage_duration_function,
do_static_initialization_or_destruction,
one_static_initialization_or_destruction,
generate_ctor_or_dtor_function): Add 'omp_target' parameter. Support
"declare target" decls. Update forward declarations.
(OMP_SSDF_IDENTIFIER): New macro.
(omp_tgt_ssdf_decls): New vec.
(get_priority_info): Initialize omp_tgt_initializations_p and
omp_tgt_destructions_p fields.
(handle_tls_init): Update call to
omp_static_initialization_or_destruction.
(c_parse_final_cleanups): Support constructors/destructors on OpenMP
offload targets.
gcc/
* omp-builtins.def (BUILT_IN_OMP_IS_INITIAL_DEVICE): New builtin.
* tree.cc (get_file_function_name): Support names for on-target
constructor/destructor functions.
libgomp/
* testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C: New
test.
* testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C: New
test.
---
gcc/cp/decl2.cc | 225 +++++++++++++++---
gcc/omp-builtins.def | 2 +
gcc/tree.cc | 6 +-
.../static-aggr-constructor-destructor-1.C | 28 +++
.../static-aggr-constructor-destructor-2.C | 31 +++
5 files changed, 257 insertions(+), 35 deletions(-)
create mode 100644
libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C
create mode 100644
libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C
diff --git a/gcc/cp/decl2.cc b/gcc/cp/decl2.cc
index f1a6df375e8..042ae4df700 100644
--- a/gcc/cp/decl2.cc
+++ b/gcc/cp/decl2.cc
@@ -65,16 +65,19 @@ typedef struct priority_info_s {
/* Nonzero if there have been any destructions at this priority
throughout the translation unit. */
int destructions_p;
+ /* Again, but specifically for OpenMP "declare target" initializations. */
+ int omp_tgt_initializations_p;
+ int omp_tgt_destructions_p;
} *priority_info;
-static tree start_objects (int, int);
+static tree start_objects (int, int, bool);
static void finish_objects (int, int, tree);
-static tree start_static_storage_duration_function (unsigned);
+static tree start_static_storage_duration_function (unsigned, bool);
static void finish_static_storage_duration_function (tree);
static priority_info get_priority_info (int);
-static void do_static_initialization_or_destruction (tree, bool);
-static void one_static_initialization_or_destruction (tree, tree, bool);
-static void generate_ctor_or_dtor_function (bool, int, location_t *);
+static void do_static_initialization_or_destruction (tree, bool, bool);
+static void one_static_initialization_or_destruction (tree, tree, bool, bool);
+static void generate_ctor_or_dtor_function (bool, int, location_t *, bool);
static int generate_ctor_and_dtor_functions_for_priority (splay_tree_node,
void *);
static tree prune_vars_needing_no_initialization (tree *);
@@ -3791,7 +3794,7 @@ generate_tls_wrapper (tree fn)
vtv_start_verification_constructor_init_function. */
static tree
-start_objects (int method_type, int initp)
+start_objects (int method_type, int initp, bool omp_target = false)
{
/* Make ctor or dtor function. METHOD_TYPE may be 'I' or 'D'. */
int module_init = 0;
@@ -3806,7 +3809,16 @@ start_objects (int method_type, int initp)
{
char type[14];
- unsigned len = sprintf (type, "sub_%c", method_type);
+ unsigned len;
+ if (omp_target)
+ /* Use "off_" signifying "offload" here. The name must be distinct
+ from the non-offload case. The format of the name is scanned in
+ tree.cc/get_file_function_name, so stick to the same length for
+ both name variants. */
+ len = sprintf (type, "off_%c", method_type);
+ else
+ len = sprintf (type, "sub_%c", method_type);
+
if (initp != DEFAULT_INIT_PRIORITY)
{
char joiner = '_';
@@ -3821,6 +3833,17 @@ start_objects (int method_type, int initp)
tree fntype = build_function_type (void_type_node, void_list_node);
tree fndecl = build_lang_decl (FUNCTION_DECL, name, fntype);
+
+ if (omp_target)
+ {
+ DECL_ATTRIBUTES (fndecl)
+ = tree_cons (get_identifier ("omp declare target"), NULL_TREE,
+ DECL_ATTRIBUTES (fndecl));
+ DECL_ATTRIBUTES (fndecl)
+ = tree_cons (get_identifier ("omp declare target nohost"), NULL_TREE,
+ DECL_ATTRIBUTES (fndecl));
+ }
+
DECL_CONTEXT (fndecl) = FROB_CONTEXT (global_namespace);
if (module_init > 0)
{
@@ -3911,6 +3934,7 @@ finish_objects (int method_type, int initp, tree body)
/* The name of the function we create to handle initializations and
destructions for objects with static storage duration. */
#define SSDF_IDENTIFIER "__static_initialization_and_destruction"
+#define OMP_SSDF_IDENTIFIER "__omp_target_static_init_and_destruction"
/* The declaration for the __INITIALIZE_P argument. */
static GTY(()) tree initialize_p_decl;
@@ -3925,6 +3949,9 @@ static GTY(()) tree ssdf_decl;
translation unit. */
static GTY(()) vec<tree, va_gc> *ssdf_decls;
+/* Same, but specifically for offloaded OpenMP "declare target" functions. */
+static GTY(()) vec<tree, va_gc> *omp_tgt_ssdf_decls;
+
/* A map from priority levels to information about that priority
level. There may be many such levels, so efficient lookup is
important. */
@@ -3943,24 +3970,37 @@ static splay_tree priority_info_map;
translation unit. */
static tree
-start_static_storage_duration_function (unsigned count)
+start_static_storage_duration_function (unsigned count, bool omp_target)
{
tree type;
tree body;
- char id[sizeof (SSDF_IDENTIFIER) + 1 /* '\0' */ + 32];
+ tree name;
- /* Create the identifier for this function. It will be of the form
- SSDF_IDENTIFIER_<number>. */
- sprintf (id, "%s_%u", SSDF_IDENTIFIER, count);
+ if (omp_target)
+ {
+ char id[sizeof (OMP_SSDF_IDENTIFIER) + 1 /* '\0' */ + 32];
+
+ /* Create the identifier for this function. It will be of the form
+ SSDF_IDENTIFIER_<number>. */
+ sprintf (id, "%s_%u", OMP_SSDF_IDENTIFIER, count);
+ name = get_identifier (id);
+ }
+ else
+ {
+ char id[sizeof (SSDF_IDENTIFIER) + 1 /* '\0' */ + 32];
+
+ /* Create the identifier for this function. It will be of the form
+ SSDF_IDENTIFIER_<number>. */
+ sprintf (id, "%s_%u", SSDF_IDENTIFIER, count);
+ name = get_identifier (id);
+ }
type = build_function_type_list (void_type_node,
integer_type_node, integer_type_node,
NULL_TREE);
/* Create the FUNCTION_DECL itself. */
- ssdf_decl = build_lang_decl (FUNCTION_DECL,
- get_identifier (id),
- type);
+ ssdf_decl = build_lang_decl (FUNCTION_DECL, name, type);
TREE_PUBLIC (ssdf_decl) = 0;
DECL_ARTIFICIAL (ssdf_decl) = 1;
@@ -3984,7 +4024,14 @@ start_static_storage_duration_function (unsigned count)
get_priority_info (DEFAULT_INIT_PRIORITY);
}
- vec_safe_push (ssdf_decls, ssdf_decl);
+ if (omp_target && !omp_tgt_ssdf_decls)
+ /* Static constructors and destructors for "declare target" variables. */
+ vec_alloc (omp_tgt_ssdf_decls, 32);
+
+ if (omp_target)
+ vec_safe_push (omp_tgt_ssdf_decls, ssdf_decl);
+ else
+ vec_safe_push (ssdf_decls, ssdf_decl);
/* Create the argument list. */
initialize_p_decl = cp_build_parm_decl
@@ -3997,6 +4044,16 @@ start_static_storage_duration_function (unsigned count)
DECL_CHAIN (initialize_p_decl) = priority_decl;
DECL_ARGUMENTS (ssdf_decl) = initialize_p_decl;
+ if (omp_target)
+ {
+ DECL_ATTRIBUTES (ssdf_decl)
+ = tree_cons (get_identifier ("omp declare target"), NULL_TREE,
+ DECL_ATTRIBUTES (ssdf_decl));
+ DECL_ATTRIBUTES (ssdf_decl)
+ = tree_cons (get_identifier ("omp declare target nohost"), NULL_TREE,
+ DECL_ATTRIBUTES (ssdf_decl));
+ }
+
/* Put the function in the global scope. */
pushdecl (ssdf_decl);
@@ -4048,6 +4105,8 @@ get_priority_info (int priority)
pi = XNEW (struct priority_info_s);
pi->initializations_p = 0;
pi->destructions_p = 0;
+ pi->omp_tgt_initializations_p = 0;
+ pi->omp_tgt_destructions_p = 0;
splay_tree_insert (priority_info_map,
(splay_tree_key) priority,
(splay_tree_value) pi);
@@ -4108,7 +4167,8 @@ fix_temporary_vars_context_r (tree *node,
are destroying it. */
static void
-one_static_initialization_or_destruction (tree decl, tree init, bool initp)
+one_static_initialization_or_destruction (tree decl, tree init, bool initp,
+ bool omp_target)
{
tree guard_if_stmt = NULL_TREE;
tree guard;
@@ -4255,7 +4315,7 @@ one_static_initialization_or_destruction (tree decl, tree
init, bool initp)
Whether initialization or destruction is performed is specified by INITP.
*/
static void
-do_static_initialization_or_destruction (tree vars, bool initp)
+do_static_initialization_or_destruction (tree vars, bool initp, bool
omp_target)
{
tree node, init_if_stmt, cond;
@@ -4298,10 +4358,14 @@ do_static_initialization_or_destruction (tree vars,
bool initp)
priority. */
priority = DECL_EFFECTIVE_INIT_PRIORITY (decl);
pi = get_priority_info (priority);
- if (initp)
+ if (initp && !omp_target)
pi->initializations_p = 1;
- else
+ else if (!omp_target)
pi->destructions_p = 1;
+ else if (initp && omp_target)
+ pi->omp_tgt_initializations_p = 1;
+ else
+ pi->omp_tgt_destructions_p = 1;
/* Conditionalize this initialization on being in the right priority
and being initializing/finalizing appropriately. */
@@ -4317,9 +4381,17 @@ do_static_initialization_or_destruction (tree vars, bool
initp)
for (; node
&& DECL_EFFECTIVE_INIT_PRIORITY (TREE_VALUE (node)) == priority;
node = TREE_CHAIN (node))
- /* Do one initialization or destruction. */
- one_static_initialization_or_destruction (TREE_VALUE (node),
- TREE_PURPOSE (node), initp);
+ {
+ tree decl = TREE_VALUE (node);
+ tree init = TREE_PURPOSE (node);
+ /* We will emit 'init' twice, and it is modified in-place during
+ gimplification. Make a copy here. */
+ if (omp_target)
+ init = copy_node (init);
+ /* Do one initialization or destruction. */
+ one_static_initialization_or_destruction (decl, init, initp,
+ omp_target);
+ }
/* Finish up the priority if-stmt body. */
finish_then_clause (priority_if_stmt);
@@ -4419,7 +4491,7 @@ write_out_vars (tree vars)
static void
generate_ctor_or_dtor_function (bool constructor_p, int priority,
- location_t *locus)
+ location_t *locus, bool omp_target)
{
input_location = *locus;
@@ -4451,13 +4523,14 @@ generate_ctor_or_dtor_function (bool constructor_p, int
priority,
arguments. */
tree fndecl;
size_t i;
- FOR_EACH_VEC_SAFE_ELT (ssdf_decls, i, fndecl)
+ vec<tree, va_gc> *walk_decls = omp_target ? omp_tgt_ssdf_decls : ssdf_decls;
+ FOR_EACH_VEC_SAFE_ELT (walk_decls, i, fndecl)
{
/* Calls to pure or const functions will expand to nothing. */
if (! (flags_from_decl_or_type (fndecl) & (ECF_CONST | ECF_PURE)))
{
if (! body)
- body = start_objects (function_key, priority);
+ body = start_objects (function_key, priority, omp_target);
tree call = cp_build_function_call_nary (fndecl, tf_warning_or_error,
build_int_cst (NULL_TREE,
@@ -4487,9 +4560,17 @@ generate_ctor_and_dtor_functions_for_priority
(splay_tree_node n, void * data)
/* Generate the functions themselves, but only if they are really
needed. */
if (pi->initializations_p)
- generate_ctor_or_dtor_function (/*constructor_p=*/true, priority, locus);
+ generate_ctor_or_dtor_function (/*constructor_p=*/true, priority, locus,
+ /*omp_target=*/false);
if (pi->destructions_p)
- generate_ctor_or_dtor_function (/*constructor_p=*/false, priority, locus);
+ generate_ctor_or_dtor_function (/*constructor_p=*/false, priority, locus,
+ /*omp_target=*/false);
+ if (pi->omp_tgt_initializations_p)
+ generate_ctor_or_dtor_function (/*constructor_p=*/true, priority, locus,
+ /*omp_target=*/true);
+ if (pi->omp_tgt_destructions_p)
+ generate_ctor_or_dtor_function (/*constructor_p=*/false, priority, locus,
+ /*omp_target=*/true);
/* Keep iterating. */
return 0;
@@ -4773,7 +4854,7 @@ handle_tls_init (void)
{
tree var = TREE_VALUE (vars);
tree init = TREE_PURPOSE (vars);
- one_static_initialization_or_destruction (var, init, true);
+ one_static_initialization_or_destruction (var, init, true, false);
/* Output init aliases even with -fno-extern-tls-init. */
if (TARGET_SUPPORTS_ALIASES && TREE_PUBLIC (var))
@@ -5087,6 +5168,7 @@ c_parse_final_cleanups (void)
int retries = 0;
unsigned ssdf_count = 0;
+ unsigned omp_target_ssdf_count = 0;
for (bool reconsider = true; reconsider; retries++)
{
reconsider = false;
@@ -5160,11 +5242,18 @@ c_parse_final_cleanups (void)
/* Set the line and file, so that it is obviously not from
the source file. */
input_location = locus_at_end_of_parsing;
- ssdf_body = start_static_storage_duration_function (ssdf_count);
+ ssdf_body
+ = start_static_storage_duration_function (ssdf_count, false);
/* First generate code to do all the initializations. */
if (vars)
- do_static_initialization_or_destruction (vars, /*initp=*/true);
+ do_static_initialization_or_destruction (vars, /*initp=*/true,
+ /*omp_target=*/false);
+
+ tree filtered_vars = NULL_TREE;
+
+ if (flag_openmp)
+ filtered_vars = copy_list (vars);
/* Then, generate code to do all the destructions. Do these
in reverse order so that the most recently constructed
@@ -5175,7 +5264,8 @@ c_parse_final_cleanups (void)
if (!flag_use_cxa_atexit && vars)
{
vars = nreverse (vars);
- do_static_initialization_or_destruction (vars, /*initp=*/false);
+ do_static_initialization_or_destruction (vars, /*initp=*/false,
+ /*omp_target=*/false);
}
else
vars = NULL_TREE;
@@ -5185,6 +5275,74 @@ c_parse_final_cleanups (void)
input_location = locus_at_end_of_parsing;
finish_static_storage_duration_function (ssdf_body);
+ if (flag_openmp)
+ {
+ /* Do all the above again for OpenMP "declare target" static
+ storage duration decls. */
+
+ /* We're only interested in "declare target" variables now. */
+ tree *fvarsp = &filtered_vars;
+ while (*fvarsp)
+ {
+ tree decl = TREE_VALUE (*fvarsp);
+
+ if (lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (decl)))
+ fvarsp = &OMP_CLAUSE_CHAIN (*fvarsp);
+ else
+ *fvarsp = OMP_CLAUSE_CHAIN (*fvarsp);
+ }
+
+ input_location = locus_at_end_of_parsing;
+ ssdf_body
+ = start_static_storage_duration_function (omp_target_ssdf_count,
+ /*omp_target=*/true);
+
+ /* As above, first generate code to do all the
+ initializations. */
+ if (filtered_vars)
+ {
+ tree nonhost_if_stmt = NULL_TREE;
+ nonhost_if_stmt = begin_if_stmt ();
+
+ /* We add an "omp declare target nohost" attribute, but (for
+ now) we still get a copy of the constructor/destructor on
+ the host. Make sure it does nothing unless we're on the
+ target device. */
+ tree fn
+ = builtin_decl_explicit (BUILT_IN_OMP_IS_INITIAL_DEVICE);
+ tree initial_dev = build_call_expr (fn, 0);
+ tree target_dev_p
+ = cp_build_binary_op (input_location, NE_EXPR, initial_dev,
+ build_int_cst (NULL_TREE, 1),
+ tf_warning_or_error);
+ finish_if_stmt_cond (target_dev_p, nonhost_if_stmt);
+
+ do_static_initialization_or_destruction (filtered_vars,
+ /*initp=*/true,
+ /*omp_target=*/true);
+ if (!flag_use_cxa_atexit && filtered_vars)
+ {
+ filtered_vars = nreverse (filtered_vars);
+ do_static_initialization_or_destruction (filtered_vars,
+ /*initp=*/false,
+ /*omp_target=*/
+ false);
+ }
+ else
+ filtered_vars = NULL_TREE;
+
+ /* Finish up nonhost if-stmt body. */
+ finish_then_clause (nonhost_if_stmt);
+ finish_if_stmt (nonhost_if_stmt);
+ }
+
+ input_location = locus_at_end_of_parsing;
+ finish_static_storage_duration_function (ssdf_body);
+
+ omp_target_ssdf_count++;
+ }
+
/* All those initializations and finalizations might cause
us to need more inline functions, more template
instantiations, etc. */
@@ -5365,7 +5523,8 @@ c_parse_final_cleanups (void)
|| module_initializer_kind ())
generate_ctor_or_dtor_function (/*constructor_p=*/true,
DEFAULT_INIT_PRIORITY,
- &locus_at_end_of_parsing);
+ &locus_at_end_of_parsing,
+ /*omp_target=*/false);
/* We're done with the splay-tree now. */
if (priority_info_map)
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index d257278b9e5..b3715b91cbb 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -68,6 +68,8 @@ DEF_GOACC_BUILTIN_ONLY (BUILT_IN_GOACC_SINGLE_COPY_START,
"GOACC_single_copy_sta
DEF_GOACC_BUILTIN_ONLY (BUILT_IN_GOACC_SINGLE_COPY_END,
"GOACC_single_copy_end",
BT_FN_VOID_PTR, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_OMP_IS_INITIAL_DEVICE, "omp_is_initial_device",
+ BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_THREAD_NUM, "omp_get_thread_num",
BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_NUM_THREADS, "omp_get_num_threads",
diff --git a/gcc/tree.cc b/gcc/tree.cc
index 13c23b67a43..aed566fcf0e 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -8769,9 +8769,11 @@ get_file_function_name (const char *type)
will be local to this file and the name is only necessary for
debugging purposes.
We also assign sub_I and sub_D sufixes to constructors called from
- the global static constructors. These are always local. */
+ the global static constructors. These are always local.
+ OpenMP "declare target" offloaded constructors/destructors use "off_I" and
+ "off_D" for the same purpose. */
else if (((type[0] == 'I' || type[0] == 'D') && targetm.have_ctors_dtors)
- || (startswith (type, "sub_")
+ || ((startswith (type, "sub_") || startswith (type, "off_"))
&& (type[4] == 'I' || type[4] == 'D')))
{
const char *file = main_input_filename;
diff --git
a/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C
b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C
new file mode 100644
index 00000000000..91d8469a150
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C
@@ -0,0 +1,28 @@
+// { dg-do run }
+
+#include <cassert>
+
+#pragma omp declare target
+
+struct str {
+ str(int x) : _x(x) { }
+ int add(str o) { return _x + o._x; }
+ int _x;
+} v1(5);
+
+#pragma omp end declare target
+
+int main()
+{
+ int res = -1;
+ str v2(2);
+
+#pragma omp target map(from:res)
+ {
+ res = v1.add(v2);
+ }
+
+ assert (res == 7);
+
+ return 0;
+}
diff --git
a/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C
b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C
new file mode 100644
index 00000000000..1bf3ee8e31c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C
@@ -0,0 +1,31 @@
+// { dg-do run }
+
+#include <cassert>
+
+#pragma omp declare target
+
+template<typename T>
+struct str {
+ str(T x) : _x(x) { }
+ T add(str o) { return _x + o._x; }
+ T _x;
+};
+
+str<long> v1(5);
+
+#pragma omp end declare target
+
+int main()
+{
+ long res = -1;
+ str<long> v2(2);
+
+#pragma omp target map(from:res)
+ {
+ res = v1.add(v2);
+ }
+
+ assert (res == 7);
+
+ return 0;
+}
--
2.29.2
