From f1bfba3a9b3f31e3e06bfd1911c9f223869ea03f Mon Sep 17 00:00:00 2001
From: Tobias Burnus <tburnus@baylibre.com>
Date: Wed, 7 Aug 2024 19:31:19 +0200
Subject: [PATCH] OpenMP: Constructors and destructors for "declare target"
static aggregates
This commit also compile-time expands (__builtin_)omp_is_initial_device for
both Fortran and C/C++ (unless, -fno-builtin-omp_is_initial_device is used).
But the main change is:
This commit 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.
Before this commit, space is allocated on the target for such aggregates,
but nothing ever constructs them properly, so they end up zero-initialised.
(See the new test static-aggr-constructor-destructor-3.C for a reason
why running constructors on the target is preferable to e.g. constructing
on the host and then copying the resulting object to the target.)
2024-08-07 Julian Brown <julian@codesourcery.com>
Tobias Burnus <tobias@baylibre.com>
gcc/ChangeLog:
* builtins.def (DEF_GOMP_BUILTIN_COMPILER): Define
DEF_GOMP_BUILTIN_COMPILER to handle the non-prefix version.
* gimple-fold.cc (gimple_fold_builtin_omp_is_initial_device): New.
(gimple_fold_builtin): Call it.
* omp-builtins.def (BUILT_IN_OMP_IS_INITIAL_DEVICE): Define.
* tree.cc (get_file_function_name): Support names for on-target
constructor/destructor functions.
gcc/cp/
* decl2.cc (tree-inline.h): Include.
(static_init_fini_fns): Bump to four entries. Update comment.
(start_objects, start_partial_init_fini_fn): Add 'omp_target'
parameter. Support "declare target" decls. Update forward declaration.
(emit_partial_init_fini_fn): Add 'host_fn' parameter. Return tree for
the created function. Support "declare target".
(OMP_SSDF_IDENTIFIER): New macro.
(partition_vars_for_init_fini): Support partitioning "declare target"
variables also.
(generate_ctor_or_dtor_function): Add 'omp_target' parameter. Support
"declare target" decls.
(c_parse_final_cleanups): Support constructors/destructors on OpenMP
offload targets.
gcc/fortran/ChangeLog:
* gfortran.h (gfc_option_t): Add disable_omp_is_initial_device.
* lang.opt (fbuiltin-): Add.
* options.cc (gfc_handle_option): Handle
-fno-builtin-omp_is_initial_device.
* f95-lang.cc (gfc_init_builtin_functions): Handle
DEF_GOMP_BUILTIN_COMPILER.
* trans-decl.cc (gfc_get_extern_function_decl): Add code to use
DEF_GOMP_BUILTIN_COMPILER for 'omp_is_initial_device'.
libgomp/ChangeLog:
* testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C: New test.
* testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C: New test.
* testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C: New test.
* testsuite/libgomp.c-c++-common/target-is-initial-host.c: New test.
* testsuite/libgomp.c-c++-common/target-is-initial-host-2.c: New test.
* testsuite/libgomp.fortran/target-is-initial-host.f: New test.
* testsuite/libgomp.fortran/target-is-initial-host.f90: New test.
* testsuite/libgomp.fortran/target-is-initial-host-2.f90: New test.
Co-authored-by: Tobias Burnus <tobias@baylibre.com>
---
gcc/builtins.def | 4 +
gcc/cp/decl2.cc | 229 +++++++++++++++---
gcc/fortran/f95-lang.cc | 9 +
gcc/fortran/gfortran.h | 1 +
gcc/fortran/lang.opt | 4 +
gcc/fortran/options.cc | 12 +
gcc/fortran/trans-decl.cc | 9 +
gcc/gimple-fold.cc | 20 ++
gcc/omp-builtins.def | 4 +
gcc/tree.cc | 6 +-
.../static-aggr-constructor-destructor-1.C | 72 ++++++
.../static-aggr-constructor-destructor-2.C | 50 ++++
.../static-aggr-constructor-destructor-3.C | 36 +++
.../target-is-initial-host-2.c | 43 ++++
.../target-is-initial-host.c | 42 ++++
.../target-is-initial-host-2.f90 | 37 +++
.../libgomp.fortran/target-is-initial-host.f | 35 +++
.../target-is-initial-host.f90 | 35 +++
18 files changed, 608 insertions(+), 40 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
create mode 100644 libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host-2.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host.c
create mode 100644 libgomp/testsuite/libgomp.fortran/target-is-initial-host-2.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/target-is-initial-host.f
create mode 100644 libgomp/testsuite/libgomp.fortran/target-is-initial-host.f90
diff --git a/gcc/builtins.def b/gcc/builtins.def
index f6f3e104f6a8..a3921aa6856a 100644
--- a/gcc/builtins.def
+++ b/gcc/builtins.def
@@ -224,6 +224,10 @@ along with GCC; see the file COPYING3. If not see
(flag_openacc \
|| flag_openmp \
|| flag_tree_parallelize_loops > 1))
+#undef DEF_GOMP_BUILTIN_COMPILER
+#define DEF_GOMP_BUILTIN_COMPILER(ENUM, NAME, TYPE, ATTRS) \
+ DEF_BUILTIN (ENUM, "__builtin_" NAME, BUILT_IN_NORMAL, TYPE, TYPE, \
+ flag_openmp, true, true, ATTRS, false, flag_openmp)
/* Builtin used by the implementation of GNU TM. These
functions are mapped to the actual implementation of the STM library. */
diff --git a/gcc/cp/decl2.cc b/gcc/cp/decl2.cc
index 695d5f8d7901..e9ae979896c5 100644
--- a/gcc/cp/decl2.cc
+++ b/gcc/cp/decl2.cc
@@ -50,6 +50,7 @@ along with GCC; see the file COPYING3. If not see
#include "asan.h"
#include "optabs-query.h"
#include "omp-general.h"
+#include "tree-inline.h"
#include "escaped_string.h"
/* Id for dumping the raw trees. */
@@ -57,14 +58,15 @@ int raw_dump_id;
extern cpp_reader *parse_in;
-static tree start_objects (bool, unsigned, bool);
+static tree start_objects (bool, unsigned, bool, bool);
static tree finish_objects (bool, unsigned, tree, bool = true);
-static tree start_partial_init_fini_fn (bool, unsigned, unsigned);
+static tree start_partial_init_fini_fn (bool, unsigned, unsigned, bool);
static void finish_partial_init_fini_fn (tree);
-static void emit_partial_init_fini_fn (bool, unsigned, tree,
- unsigned, location_t);
+static tree emit_partial_init_fini_fn (bool, unsigned, tree,
+ unsigned, location_t, tree);
static void one_static_initialization_or_destruction (bool, tree, tree);
-static void generate_ctor_or_dtor_function (bool, unsigned, tree, location_t);
+static void generate_ctor_or_dtor_function (bool, unsigned, tree, location_t,
+ bool);
static tree prune_vars_needing_no_initialization (tree *);
static void write_out_vars (tree);
static void import_export_class (tree);
@@ -166,9 +168,10 @@ struct priority_map_traits
typedef hash_map<unsigned/*Priority*/, tree/*List*/,
priority_map_traits> priority_map_t;
-/* A pair of such hash tables, indexed by initp -- one for fini and
- one for init. The fini table is only ever used when !cxa_atexit. */
-static GTY(()) priority_map_t *static_init_fini_fns[2];
+/* Two pairs of such hash tables, for the host and an OpenMP offload device.
+ Each pair has one priority map for fini and one for init. The fini tables
+ are only ever used when !cxa_atexit. */
+static GTY(()) priority_map_t *static_init_fini_fns[4];
/* Nonzero if we're done parsing and into end-of-file activities.
2 if all templates have been instantiated.
@@ -4048,7 +4051,8 @@ generate_tls_wrapper (tree fn)
/* Start a global constructor or destructor function. */
static tree
-start_objects (bool initp, unsigned priority, bool has_body)
+start_objects (bool initp, unsigned priority, bool has_body,
+ bool omp_target = false)
{
bool default_init = initp && priority == DEFAULT_INIT_PRIORITY;
bool is_module_init = default_init && module_global_init_needed ();
@@ -4062,7 +4066,15 @@ start_objects (bool initp, unsigned priority, bool has_body)
/* We use `I' to indicate initialization and `D' to indicate
destruction. */
- unsigned len = sprintf (type, "sub_%c", initp ? 'I' : 'D');
+ 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", initp ? 'I' : 'D');
+ else
+ len = sprintf (type, "sub_%c", initp ? 'I' : 'D');
if (priority != DEFAULT_INIT_PRIORITY)
{
char joiner = '_';
@@ -4077,6 +4089,17 @@ start_objects (bool initp, unsigned priority, bool has_body)
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 (is_module_init)
{
@@ -4161,34 +4184,53 @@ finish_objects (bool initp, unsigned priority, tree body, bool startp)
/* 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"
/* Begins the generation of the function that will handle all
initialization or destruction of objects with static storage
duration at PRIORITY.
- It is assumed that this function will only be called once. */
+ It is assumed that this function will be called once for the host, and once
+ for an OpenMP offload target. */
static tree
-start_partial_init_fini_fn (bool initp, unsigned priority, unsigned count)
+start_partial_init_fini_fn (bool initp, unsigned priority, unsigned count,
+ bool omp_target)
{
- char id[sizeof (SSDF_IDENTIFIER) + 1 /* '\0' */ + 32];
+ char id[MAX (sizeof (SSDF_IDENTIFIER), sizeof (OMP_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);
-
+ SSDF_IDENTIFIER_<number> if not omp_target and otherwise
+ OMP_SSDF_IDENTIFIER_<number>. */
+ sprintf (id, "%s_%u", omp_target ? OMP_SSDF_IDENTIFIER : SSDF_IDENTIFIER,
+ count);
+ name = get_identifier (id);
tree type = build_function_type (void_type_node, void_list_node);
/* Create the FUNCTION_DECL itself. */
- tree fn = build_lang_decl (FUNCTION_DECL, get_identifier (id), type);
+ tree fn = build_lang_decl (FUNCTION_DECL, name, type);
TREE_PUBLIC (fn) = 0;
DECL_ARTIFICIAL (fn) = 1;
+ if (omp_target)
+ {
+ DECL_ATTRIBUTES (fn)
+ = tree_cons (get_identifier ("omp declare target"), NULL_TREE,
+ DECL_ATTRIBUTES (fn));
+ DECL_ATTRIBUTES (fn)
+ = tree_cons (get_identifier ("omp declare target nohost"), NULL_TREE,
+ DECL_ATTRIBUTES (fn));
+ }
+
+ int idx = initp + 2 * omp_target;
+
/* Put this function in the list of functions to be called from the
static constructors and destructors. */
- if (!static_init_fini_fns[initp])
- static_init_fini_fns[initp] = priority_map_t::create_ggc ();
- auto &slot = static_init_fini_fns[initp]->get_or_insert (priority);
+ if (!static_init_fini_fns[idx])
+ static_init_fini_fns[idx] = priority_map_t::create_ggc ();
+ auto &slot = static_init_fini_fns[idx]->get_or_insert (priority);
slot = tree_cons (fn, NULL_TREE, slot);
/* Put the function in the global scope. */
@@ -4384,22 +4426,74 @@ one_static_initialization_or_destruction (bool initp, tree decl, tree init)
a TREE_LIST of VAR_DECL with static storage duration.
Whether initialization or destruction is performed is specified by INITP. */
-static void
+static tree
emit_partial_init_fini_fn (bool initp, unsigned priority, tree vars,
- unsigned counter, location_t locus)
+ unsigned counter, location_t locus, tree host_fn)
{
input_location = locus;
- tree body = start_partial_init_fini_fn (initp, priority, counter);
+ bool omp_target = (host_fn != NULL_TREE);
+ tree body = start_partial_init_fini_fn (initp, priority, counter, omp_target);
+ tree fndecl = current_function_decl;
+
+ tree nonhost_if_stmt = NULL_TREE;
+ if (omp_target)
+ {
+ 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,
+ integer_one_node, tf_warning_or_error);
+ finish_if_stmt_cond (target_dev_p, nonhost_if_stmt);
+ }
for (tree node = vars; node; node = TREE_CHAIN (node))
- /* Do one initialization or destruction. */
- one_static_initialization_or_destruction (initp, TREE_VALUE (node),
- TREE_PURPOSE (node));
+ {
+ 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)
+ {
+ /* We've already emitted INIT in the host version of the ctor/dtor
+ function. We need to deep-copy it (including new versions of
+ local variables introduced, etc.) for use in the target
+ ctor/dtor function. */
+ copy_body_data id;
+ hash_map<tree, tree> decl_map;
+ memset (&id, 0, sizeof (id));
+ id.src_fn = host_fn;
+ id.dst_fn = current_function_decl;
+ id.src_cfun = DECL_STRUCT_FUNCTION (id.src_fn);
+ id.decl_map = &decl_map;
+ id.copy_decl = copy_decl_no_change;
+ id.transform_call_graph_edges = CB_CGE_DUPLICATE;
+ id.transform_new_cfg = true;
+ id.transform_return_to_modify = false;
+ id.eh_lp_nr = 0;
+ walk_tree (&init, copy_tree_body_r, &id, NULL);
+ }
+ /* Do one initialization or destruction. */
+ one_static_initialization_or_destruction (initp, decl, init);
+ }
+
+ if (omp_target)
+ {
+ /* Finish up nonhost if-stmt body. */
+ finish_then_clause (nonhost_if_stmt);
+ finish_if_stmt (nonhost_if_stmt);
+ }
/* Finish up the static storage duration function for this
round. */
input_location = locus;
finish_partial_init_fini_fn (body);
+
+ return fndecl;
}
/* VARS is a list of variables with static storage duration which may
@@ -4462,7 +4556,7 @@ prune_vars_needing_no_initialization (tree *vars)
This reverses the variable ordering. */
void
-partition_vars_for_init_fini (tree var_list, priority_map_t *(&parts)[2])
+partition_vars_for_init_fini (tree var_list, priority_map_t *(&parts)[4])
{
for (auto node = var_list; node; node = TREE_CHAIN (node))
{
@@ -4488,6 +4582,30 @@ partition_vars_for_init_fini (tree var_list, priority_map_t *(&parts)[2])
auto &slot = parts[false]->get_or_insert (priority);
slot = tree_cons (NULL_TREE, decl, slot);
}
+
+ if (flag_openmp
+ && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)))
+ {
+ priority_map_t **omp_parts = parts + 2;
+
+ if (init || (flag_use_cxa_atexit && has_cleanup))
+ {
+ // Add to initialization list.
+ if (!omp_parts[true])
+ omp_parts[true] = priority_map_t::create_ggc ();
+ auto &slot = omp_parts[true]->get_or_insert (priority);
+ slot = tree_cons (init, decl, slot);
+ }
+
+ if (!flag_use_cxa_atexit && has_cleanup)
+ {
+ // Add to finalization list.
+ if (!omp_parts[false])
+ omp_parts[false] = priority_map_t::create_ggc ();
+ auto &slot = omp_parts[false]->get_or_insert (priority);
+ slot = tree_cons (NULL_TREE, decl, slot);
+ }
+ }
}
}
@@ -4515,10 +4633,10 @@ write_out_vars (tree vars)
static void
generate_ctor_or_dtor_function (bool initp, unsigned priority,
- tree fns, location_t locus)
+ tree fns, location_t locus, bool omp_target)
{
input_location = locus;
- tree body = start_objects (initp, priority, bool (fns));
+ tree body = start_objects (initp, priority, bool (fns), omp_target);
if (fns)
{
@@ -5190,7 +5308,7 @@ c_parse_final_cleanups (void)
auto_vec<tree> consteval_vtables;
int retries = 0;
- unsigned ssdf_count = 0;
+ unsigned ssdf_count = 0, omp_ssdf_count = 0;
for (bool reconsider = true; reconsider; retries++)
{
reconsider = false;
@@ -5253,8 +5371,9 @@ c_parse_final_cleanups (void)
write_out_vars (vars);
function_depth++; // Disable GC
- priority_map_t *parts[2] = {nullptr, nullptr};
+ priority_map_t *parts[4] = {nullptr, nullptr, nullptr, nullptr};
partition_vars_for_init_fini (vars, parts);
+ tree host_init_fini[2] = { NULL_TREE, NULL_TREE };
for (unsigned initp = 2; initp--;)
if (parts[initp])
@@ -5265,10 +5384,32 @@ c_parse_final_cleanups (void)
// Partitioning kept the vars in reverse order.
// We only want that for dtors.
list = nreverse (list);
- emit_partial_init_fini_fn (initp, iter.first, list,
- ssdf_count++,
- locus_at_end_of_parsing);
+ host_init_fini[initp]
+ = emit_partial_init_fini_fn (initp, iter.first, list,
+ ssdf_count++,
+ locus_at_end_of_parsing,
+ NULL_TREE);
}
+
+ if (flag_openmp)
+ {
+ priority_map_t **omp_parts = parts + 2;
+ for (unsigned initp = 2; initp--;)
+ if (omp_parts[initp])
+ for (auto iter : *omp_parts[initp])
+ {
+ auto list = iter.second;
+ if (initp)
+ // Partitioning kept the vars in reverse order.
+ // We only want that for dtors.
+ list = nreverse (list);
+ emit_partial_init_fini_fn (initp, iter.first, list,
+ omp_ssdf_count++,
+ locus_at_end_of_parsing,
+ host_init_fini[initp]);
+ }
+ }
+
function_depth--; // Re-enable GC
/* All those initializations and finalizations might cause
@@ -5439,6 +5580,10 @@ c_parse_final_cleanups (void)
for (auto iter : *static_init_fini_fns[true])
iter.second = nreverse (iter.second);
+ if (flag_openmp && static_init_fini_fns[2 + true])
+ for (auto iter : *static_init_fini_fns[2 + true])
+ iter.second = nreverse (iter.second);
+
/* Now we've instantiated all templates. Now we can escalate the functions
we squirreled away earlier. */
process_and_check_pending_immediate_escalating_fns ();
@@ -5457,7 +5602,7 @@ c_parse_final_cleanups (void)
{
input_location = locus_at_end_of_parsing;
tree body = start_partial_init_fini_fn (true, DEFAULT_INIT_PRIORITY,
- ssdf_count++);
+ ssdf_count++, false);
/* For Objective-C++, we may need to initialize metadata found
in this module. This must be done _before_ any other static
initializations. */
@@ -5476,18 +5621,26 @@ c_parse_final_cleanups (void)
static_init_fini_fns[true] = priority_map_t::create_ggc ();
if (static_init_fini_fns[true]->get_or_insert (DEFAULT_INIT_PRIORITY))
has_module_inits = true;
+
+ if (flag_openmp)
+ {
+ if (!static_init_fini_fns[2 + true])
+ static_init_fini_fns[2 + true] = priority_map_t::create_ggc ();
+ static_init_fini_fns[2 + true]->get_or_insert (DEFAULT_INIT_PRIORITY);
+ }
}
/* Generate initialization and destruction functions for all
priorities for which they are required. They have C-language
linkage. */
push_lang_context (lang_name_c);
- for (unsigned initp = 2; initp--;)
+ for (unsigned initp = 4; initp--;)
if (static_init_fini_fns[initp])
{
for (auto iter : *static_init_fini_fns[initp])
- generate_ctor_or_dtor_function (initp, iter.first, iter.second,
- locus_at_end_of_parsing);
+ generate_ctor_or_dtor_function (initp & 1, iter.first, iter.second,
+ locus_at_end_of_parsing,
+ (initp & 2) != 0);
static_init_fini_fns[initp] = nullptr;
}
pop_lang_context ();
diff --git a/gcc/fortran/f95-lang.cc b/gcc/fortran/f95-lang.cc
index 67fda27aa3ed..770f31b221d6 100644
--- a/gcc/fortran/f95-lang.cc
+++ b/gcc/fortran/f95-lang.cc
@@ -1271,10 +1271,13 @@ gfc_init_builtin_functions (void)
attr);
#undef DEF_GOMP_BUILTIN
#define DEF_GOMP_BUILTIN(code, name, type, attr) /* ignore */
+#undef DEF_GOMP_BUILTIN_COMPILER
+#define DEF_GOMP_BUILTIN_COMPILER(code, name, type, attr) /* ignore */
#include "../omp-builtins.def"
#undef DEF_GOACC_BUILTIN
#undef DEF_GOACC_BUILTIN_COMPILER
#undef DEF_GOMP_BUILTIN
+#undef DEF_GOMP_BUILTIN_COMPILER
}
if (flag_openmp || flag_openmp_simd || flag_tree_parallelize_loops)
@@ -1287,10 +1290,16 @@ gfc_init_builtin_functions (void)
#define DEF_GOMP_BUILTIN(code, name, type, attr) \
gfc_define_builtin ("__builtin_" name, builtin_types[type], \
code, name, attr);
+#undef DEF_GOMP_BUILTIN_COMPILER
+#define DEF_GOMP_BUILTIN_COMPILER(code, name, type, attr) \
+ if (flag_openmp) \
+ gfc_define_builtin ("__builtin_" name, builtin_types[type], \
+ code, name, attr);
#include "../omp-builtins.def"
#undef DEF_GOACC_BUILTIN
#undef DEF_GOACC_BUILTIN_COMPILER
#undef DEF_GOMP_BUILTIN
+#undef DEF_GOMP_BUILTIN_COMPILER
tree gomp_alloc = builtin_decl_explicit (BUILT_IN_GOMP_ALLOC);
tree two = build_int_cst (integer_type_node, 2);
DECL_ATTRIBUTES (gomp_alloc)
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 6207ad6ed1df..8d89797412e8 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -3161,6 +3161,7 @@ typedef struct
int flag_init_logical;
int flag_init_character;
char flag_init_character_value;
+ int disable_omp_is_initial_device;
int fpe;
int fpe_summary;
diff --git a/gcc/fortran/lang.opt b/gcc/fortran/lang.opt
index 5efd4a0129a6..5cf7b4922540 100644
--- a/gcc/fortran/lang.opt
+++ b/gcc/fortran/lang.opt
@@ -420,6 +420,10 @@ Fortran RejectNegative Joined UInteger Var(flag_blas_matmul_limit) Init(30)
fbuilding-libgfortran
Fortran Undocumented Var(flag_building_libgfortran)
+fbuiltin-
+Fortran Joined
+; Documented in C
+
fcheck-array-temporaries
Fortran
Produce a warning at runtime if a array temporary has been created for a procedure argument.
diff --git a/gcc/fortran/options.cc b/gcc/fortran/options.cc
index 9bc5ef3779d3..d8c5c8e62fc7 100644
--- a/gcc/fortran/options.cc
+++ b/gcc/fortran/options.cc
@@ -843,6 +843,18 @@ gfc_handle_option (size_t scode, const char *arg, HOST_WIDE_INT value,
/* Set (or unset) the DEC extension flags. */
set_dec_flags (value);
break;
+
+ case OPT_fbuiltin_:
+ /* We only handle -fno-builtin-omp_is_initial_device. */
+ if (value)
+ return false; /* Not supported. */
+ if (!strcmp ("omp_is_initial_device", arg))
+ gfc_option.disable_omp_is_initial_device = true;
+ else
+ warning (0, "command-line option %<-fno-builtin-%s%> is not valid for "
+ "Fortran", arg);
+ break;
+
}
Fortran_handle_option_auto (&global_options, &global_options_set,
diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc
index 0fdc41b1784b..ca6a515a1800 100644
--- a/gcc/fortran/trans-decl.cc
+++ b/gcc/fortran/trans-decl.cc
@@ -2214,6 +2214,15 @@ gfc_get_extern_function_decl (gfc_symbol * sym, gfc_actual_arglist *actual_args,
to know that. */
gcc_assert (!(sym->attr.entry || sym->attr.entry_master));
+ if (!gfc_option.disable_omp_is_initial_device
+ && flag_openmp && sym->attr.function && sym->ts.type == BT_LOGICAL
+ && !strcmp (sym->name, "omp_is_initial_device"))
+ {
+ sym->backend_decl
+ = builtin_decl_explicit (BUILT_IN_OMP_IS_INITIAL_DEVICE);
+ return sym->backend_decl;
+ }
+
if (sym->attr.proc_pointer)
return get_proc_pointer_decl (sym);
diff --git a/gcc/gimple-fold.cc b/gcc/gimple-fold.cc
index c20102f73f59..18d7a6b176db 100644
--- a/gcc/gimple-fold.cc
+++ b/gcc/gimple-fold.cc
@@ -4043,6 +4043,23 @@ gimple_fold_builtin_strlen (gimple_stmt_iterator *gsi)
return false;
}
+static bool
+gimple_fold_builtin_omp_is_initial_device (gimple_stmt_iterator *gsi)
+{
+#if ACCEL_COMPILER
+ replace_call_with_value (gsi, integer_zero_node);
+ return true;
+#else
+ if (!ENABLE_OFFLOADING || symtab->state == EXPANSION)
+ {
+ replace_call_with_value (gsi, integer_one_node);
+ return true;
+ }
+#endif
+ return false;
+}
+
+
/* Fold a call to __builtin_acc_on_device. */
static bool
@@ -5220,6 +5237,9 @@ gimple_fold_builtin (gimple_stmt_iterator *gsi)
case BUILT_IN_ACC_ON_DEVICE:
return gimple_fold_builtin_acc_on_device (gsi,
gimple_call_arg (stmt, 0));
+ case BUILT_IN_OMP_IS_INITIAL_DEVICE:
+ return gimple_fold_builtin_omp_is_initial_device (gsi);
+
case BUILT_IN_REALLOC:
return gimple_fold_builtin_realloc (gsi);
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index 044d5d087b6d..7b49ef1c0e53 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -23,6 +23,7 @@ along with GCC; see the file COPYING3. If not see
DEF_GOACC_BUILTIN (ENUM, NAME, TYPE, ATTRS)
DEF_GOACC_BUILTIN_COMPILER (ENUM, NAME, TYPE, ATTRS)
DEF_GOMP_BUILTIN (ENUM, NAME, TYPE, ATTRS)
+ DEF_GOMP_BUILTIN_COMPILER (ENUM, NAME, TYPE, ATTRS)
See builtins.def for details. */
@@ -68,6 +69,9 @@ 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_COMPILER (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 a2d431662bd5..17a5cea7c252 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -8908,9 +8908,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 000000000000..403a071c0c01
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C
@@ -0,0 +1,72 @@
+// { dg-do run }
+// { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" }
+// { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } }
+
+// { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } }
+// { dg-final { scan-tree-dump-times "_GLOBAL__off_I_v1" 1 "gimple" } }
+// { dg-final { scan-tree-dump-times "__omp_target_static_init_and_destruction" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "__attribute__\\(\\(\[^\n\r]*omp declare target nohost" 2 "gimple" } }
+
+// { dg-final { scan-tree-dump-not "omp_is_initial_device" "optimized" } }
+// { dg-final { scan-tree-dump-not "__omp_target_static_init_and_destruction" "optimized" } }
+// FIXME: should be '-not' not '-times' 1:
+// { dg-final { scan-tree-dump-times "void _GLOBAL__off_I_v1" 1 "optimized" } }
+// { dg-final { scan-tree-dump-times "__attribute__\\(\\(\[^\n\r]*omp declare target nohost" 1 "optimized" } }
+
+// { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-not "omp_initial_device;" "optimized" { target offload_target_amdgcn } } }
+// { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "v1\\._x = 5;" "optimized" { target offload_target_amdgcn } } }
+// { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-not "omp_initial_device;" "optimized" { target offload_target_nvptx } } }
+// { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "v1\\._x = 5;" "optimized" { target offload_target_nvptx } } }
+
+#include <cassert>
+#include <omp.h>
+
+#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
+
+void check_host()
+{
+ assert (v1._x == 5);
+}
+
+void check_devs()
+{
+ for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++)
+ {
+ int res = 99, dev_num = 98;
+ #pragma omp target map(from: res, dev_num) device(dev)
+ {
+ res = v1._x;
+ dev_num = omp_get_device_num();
+ }
+ assert (res == 5);
+ if (dev == omp_initial_device)
+ assert (dev_num == omp_get_num_devices());
+ else
+ assert (dev_num == dev);
+ }
+}
+
+int main()
+{
+ int res = -1;
+ str v2(2);
+
+#pragma omp target map(from:res)
+ {
+ res = v1.add(v2);
+ }
+
+ assert (res == 7);
+ check_host();
+ check_devs();
+
+ 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 000000000000..6dd4260a522c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C
@@ -0,0 +1,50 @@
+// { dg-do run }
+// { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" }
+// { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } }
+
+// { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } }
+// { dg-final { scan-tree-dump-times "_GLOBAL__off_I_v1" 1 "gimple" } }
+// { dg-final { scan-tree-dump-times "__omp_target_static_init_and_destruction" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "__attribute__\\(\\(\[^\n\r]*omp declare target nohost" 2 "gimple" } }
+
+// { dg-final { scan-tree-dump-not "omp_is_initial_device" "optimized" } }
+// { dg-final { scan-tree-dump-not "__omp_target_static_init_and_destruction" "optimized" } }
+// FIXME: should be '-not' not '-times' 1:
+// { dg-final { scan-tree-dump-times "void _GLOBAL__off_I_" 1 "optimized" } }
+// { dg-final { scan-tree-dump-times "__attribute__\\(\\(\[^\n\r]*omp declare target nohost" 1 "optimized" } }
+
+// { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-not "omp_initial_device;" "optimized" { target offload_target_amdgcn } } }
+// { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "v1\\._x = 5;" "optimized" { target offload_target_amdgcn } } }
+// { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-not "omp_initial_device;" "optimized" { target offload_target_nvptx } } }
+// { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "v1\\._x = 5;" "optimized" { target offload_target_nvptx } } }
+
+
+#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;
+}
diff --git a/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C
new file mode 100644
index 000000000000..8d4aff21cd79
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C
@@ -0,0 +1,36 @@
+// { dg-do run }
+
+#include <cassert>
+
+#pragma omp declare target
+
+struct item {
+ item(item *p, int v) : prev(p), val(v) { }
+ int get() { return prev ? prev->get() * val : val; }
+ item *prev;
+ int val;
+};
+
+/* This case demonstrates why constructing on the host and then copying to
+ the target would be less desirable. With on-target construction, "prev"
+ for each 'item' will be a device pointer, not a host pointer. */
+item hubert1(nullptr, 3);
+item hubert2(&hubert1, 5);
+item hubert3(&hubert2, 7);
+item hubert4(&hubert3, 11);
+
+#pragma omp end declare target
+
+int main()
+{
+ int res = -1;
+
+#pragma omp target map(from:res)
+ {
+ res = hubert4.get ();
+ }
+
+ assert (res == 1155);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host-2.c
new file mode 100644
index 000000000000..313d188a871a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host-2.c
@@ -0,0 +1,43 @@
+/* { dg-do run } */
+/* { dg-additional-options "-fno-builtin-omp_is_initial_device" } */
+
+/* Check whether 'omp_is_initial_device()' is NOT compile-time optimized. */
+
+/* { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" } */
+/* { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } } */
+
+/* { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } } */
+
+/* { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "optimized" } } */
+
+/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-times "omp_is_initial_device" 1 "optimized" { target offload_target_amdgcn } } } */
+/* { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-times "omp_is_initial_device" 1 "optimized" { target offload_target_nvptx } } } */
+
+
+#include <omp.h>
+
+int
+main ()
+{
+ int is_initial, dev_num, initial;
+ initial = omp_get_initial_device();
+ for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++)
+ {
+ is_initial = dev_num = 99;
+ #pragma omp target map(from: is_initial, dev_num) device(dev)
+ {
+ is_initial = omp_is_initial_device ();
+ dev_num = omp_get_device_num ();
+ }
+ if (dev == omp_initial_device || dev == initial)
+ {
+ if (dev_num != initial || is_initial != 1)
+ __builtin_abort ();
+ }
+ else
+ {
+ if (dev_num != dev || is_initial != 0)
+ __builtin_abort ();
+ }
+ }
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host.c b/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host.c
new file mode 100644
index 000000000000..423727ce55f0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host.c
@@ -0,0 +1,42 @@
+/* { dg-do run } */
+
+/* Check whether 'omp_is_initial_device()' is properly compile-time optimized. */
+
+/* { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" } */
+/* { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } } */
+
+/* { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } } */
+
+/* { dg-final { scan-tree-dump-not "omp_is_initial_device" "optimized" } } */
+
+/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_amdgcn } } } */
+/* { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_nvptx } } } */
+
+
+#include <omp.h>
+
+int
+main ()
+{
+ int is_initial, dev_num, initial;
+ initial = omp_get_initial_device();
+ for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++)
+ {
+ is_initial = dev_num = 99;
+ #pragma omp target map(from: is_initial, dev_num) device(dev)
+ {
+ is_initial = omp_is_initial_device ();
+ dev_num = omp_get_device_num ();
+ }
+ if (dev == omp_initial_device || dev == initial)
+ {
+ if (dev_num != initial || is_initial != 1)
+ __builtin_abort ();
+ }
+ else
+ {
+ if (dev_num != dev || is_initial != 0)
+ __builtin_abort ();
+ }
+ }
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-is-initial-host-2.f90 b/libgomp/testsuite/libgomp.fortran/target-is-initial-host-2.f90
new file mode 100644
index 000000000000..e06ced206452
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-is-initial-host-2.f90
@@ -0,0 +1,37 @@
+! { dg-additional-options "-fno-builtin-omp_is_initial_device" }
+
+! Check whether 'omp_is_initial_device()' is NOT compile-time optimized. */
+
+! { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" }
+! { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } }
+
+! { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } }
+
+! { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "optimized" } }
+
+! { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-times "omp_is_initial_device" 1 "optimized" { target offload_target_amdgcn } } }
+! { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-times "omp_is_initial_device" 1 "optimized" { target offload_target_nvptx } } }
+
+
+program main
+ use omp_lib
+ implicit none (type, external)
+ integer :: dev_num, initial, dev
+ logical :: is_initial
+
+ initial = omp_get_initial_device()
+ do dev = omp_initial_device, omp_get_num_devices()
+ dev_num = 99
+ !$omp target map(from: is_initial, dev_num) device(dev)
+ is_initial = omp_is_initial_device ()
+ dev_num = omp_get_device_num ()
+ !$omp end target
+ if (dev == omp_initial_device .or. dev == initial) then
+ if (dev_num /= initial .or. .not. is_initial) &
+ stop 1
+ else
+ if (dev_num /= dev .or. is_initial) &
+ stop 2
+ end if
+ end do
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f b/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f
new file mode 100644
index 000000000000..fec4a3f1c1ce
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f
@@ -0,0 +1,35 @@
+! Check whether 'omp_is_initial_device()' is properly compile-time optimized. */
+
+! { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" }
+! { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } }
+
+! { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } }
+
+! { dg-final { scan-tree-dump-not "omp_is_initial_device" "optimized" } }
+
+! { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_amdgcn } } }
+! { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_nvptx } } }
+
+
+ program main
+ implicit none (type, external)
+ include "omp_lib.h"
+ integer :: dev_num, initial, dev
+ logical :: is_initial
+
+ initial = omp_get_initial_device()
+ do dev = omp_initial_device, omp_get_num_devices()
+ dev_num = 99
+!$omp target map(from: is_initial, dev_num) device(dev)
+ is_initial = omp_is_initial_device ()
+ dev_num = omp_get_device_num ()
+!$omp end target
+ if (dev == omp_initial_device .or. dev == initial) then
+ if (dev_num /= initial .or. .not. is_initial) &
+ & stop 1
+ else
+ if (dev_num /= dev .or. is_initial) &
+ & stop 2
+ end if
+ end do
+ end
diff --git a/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f90 b/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f90
new file mode 100644
index 000000000000..f8a645fc488f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f90
@@ -0,0 +1,35 @@
+! Check whether 'omp_is_initial_device()' is properly compile-time optimized. */
+
+! { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" }
+! { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } }
+
+! { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } }
+
+! { dg-final { scan-tree-dump-not "omp_is_initial_device" "optimized" } }
+
+! { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_amdgcn } } }
+! { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_nvptx } } }
+
+
+program main
+ use omp_lib
+ implicit none (type, external)
+ integer :: dev_num, initial, dev
+ logical :: is_initial
+
+ initial = omp_get_initial_device()
+ do dev = omp_initial_device, omp_get_num_devices()
+ dev_num = 99
+ !$omp target map(from: is_initial, dev_num) device(dev)
+ is_initial = omp_is_initial_device ()
+ dev_num = omp_get_device_num ()
+ !$omp end target
+ if (dev == omp_initial_device .or. dev == initial) then
+ if (dev_num /= initial .or. .not. is_initial) &
+ stop 1
+ else
+ if (dev_num /= dev .or. is_initial) &
+ stop 2
+ end if
+ end do
+end
--
GitLab