diff --git a/gcc/builtins.def b/gcc/builtins.def
index f6f3e104f6a8219eff4ac26fda588da808e4b7c9..a3921aa6856a72c9eaa28e772be6e138a6bf8503 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 695d5f8d7901aeda8b75880df7bdfca8b1c84ac4..e9ae979896c5770e2e2ec54e2e99278df18ae119 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 67fda27aa3edc346bf3bc3f69bddf9f93c55d13d..770f31b221d6bf841ff79d09f2c6d76588698a79 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 6207ad6ed1dfdb9824103495a48824bc29804631..8d89797412e875fe8f4c8d65cd0f22e70df90b57 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 5efd4a0129a6ee426e473d93c57d049bbcfb6ed3..5cf7b49225400a35de2e115b24e89a8761f0f01c 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 9bc5ef3779d398625327384eee479d7fd79f1d91..d8c5c8e62fc754be3e24046e68438ad17ed02ee8 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 0fdc41b1784b91701775ee1acabddb24a7236f8d..ca6a515a1800a5760e128dba90d8ffbd29f6c0f5 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 c20102f73f59758fab313cd3bffeb3068ce1b012..18d7a6b176db7ee2d7cd68dcf418b596bbf0dee0 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 044d5d087b6d32b1b5625f8772986e9a0923bdce..7b49ef1c0e534b67e9d7a9f6979ae7e733b5940a 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 a2d431662bd5dc1124e37f2a3b7fe42d3a34d0d7..17a5cea7c252b68d6b75d2be9a278075214858cb 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 0000000000000000000000000000000000000000..403a071c0c01c4bf187c85ccc5b249b348345892
--- /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 0000000000000000000000000000000000000000..6dd4260a522cb9783be0d5d18f42c492344d4121
--- /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 0000000000000000000000000000000000000000..8d4aff21cd793a010eabeeba8e8fc533459e840f
--- /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 0000000000000000000000000000000000000000..313d188a871a7064a1687f6a005d256ce920ae6b
--- /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 0000000000000000000000000000000000000000..423727ce55f0c611ca0e13d6a8fb36358a056c36
--- /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 0000000000000000000000000000000000000000..e06ced2064522fde734f0469c794fdcc0d50ce59
--- /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 0000000000000000000000000000000000000000..fec4a3f1c1cebaf4b02a9916713564b23768edf3
--- /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 0000000000000000000000000000000000000000..f8a645fc488fc166b294e256ea9c3a172c31d7c2
--- /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