diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 97e3b23b5d20da56b362f64b6af4f4df67d36feb..9c02141e2c66fdfc2e27db2c3b4f34f074990c79 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -20915,6 +20915,10 @@ c_parser_omp_teams (location_t loc, c_parser *parser,
static tree
c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
{
+ if (flag_openmp)
+ omp_requires_mask
+ = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
tree clauses
= c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
"#pragma omp target data");
@@ -21010,6 +21014,10 @@ c_parser_omp_target_update (location_t loc, c_parser *parser,
return false;
}
+ if (flag_openmp)
+ omp_requires_mask
+ = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
tree stmt = make_node (OMP_TARGET_UPDATE);
TREE_TYPE (stmt) = void_type_node;
OMP_TARGET_UPDATE_CLAUSES (stmt) = clauses;
@@ -21057,6 +21065,10 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
return true;
}
+ if (flag_openmp)
+ omp_requires_mask
+ = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
tree clauses
= c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
"#pragma omp target enter data");
@@ -21151,6 +21163,10 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
return true;
}
+ if (flag_openmp)
+ omp_requires_mask
+ = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
tree clauses
= c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
"#pragma omp target exit data");
@@ -22779,9 +22795,6 @@ c_parser_omp_requires (c_parser *parser)
c_parser_skip_to_pragma_eol (parser, false);
return;
}
- if (p && this_req != OMP_REQUIRES_DYNAMIC_ALLOCATORS)
- sorry_at (cloc, "%qs clause on %<requires%> directive not "
- "supported yet", p);
if (p)
c_parser_consume_token (parser);
if (this_req)
diff --git a/gcc/config/gcn/mkoffload.cc b/gcc/config/gcn/mkoffload.cc
index ed93ae844e47530eeeb45306191c376538e9b91c..b8b3fecfcb4176c78c7f9bf94203eb3f4905fd9f 100644
--- a/gcc/config/gcn/mkoffload.cc
+++ b/gcc/config/gcn/mkoffload.cc
@@ -611,6 +611,7 @@ process_asm (FILE *in, FILE *out, FILE *cfile)
struct regcount *regcounts = XOBFINISH (®counts_os, struct regcount *);
fprintf (cfile, "#include <stdlib.h>\n");
+ fprintf (cfile, "#include <stdint.h>\n");
fprintf (cfile, "#include <stdbool.h>\n\n");
fprintf (cfile, "static const int gcn_num_vars = %d;\n\n", var_count);
@@ -664,7 +665,7 @@ process_asm (FILE *in, FILE *out, FILE *cfile)
/* Embed an object file into a C source file. */
static void
-process_obj (FILE *in, FILE *cfile)
+process_obj (FILE *in, FILE *cfile, uint32_t omp_requires)
{
size_t len = 0;
const char *input = read_file (in, &len);
@@ -692,16 +693,18 @@ process_obj (FILE *in, FILE *cfile)
fprintf (cfile,
"static const struct gcn_image_desc {\n"
+ " uintptr_t omp_requires_mask;\n"
" const struct gcn_image *gcn_image;\n"
" unsigned kernel_count;\n"
" const struct hsa_kernel_description *kernel_infos;\n"
" unsigned global_variable_count;\n"
"} target_data = {\n"
+ " %d,\n"
" &gcn_image,\n"
" sizeof (gcn_kernels) / sizeof (gcn_kernels[0]),\n"
" gcn_kernels,\n"
" gcn_num_vars\n"
- "};\n\n");
+ "};\n\n", omp_requires);
fprintf (cfile,
"#ifdef __cplusplus\n"
@@ -1077,9 +1080,27 @@ main (int argc, char **argv)
unsetenv ("COMPILER_PATH");
unsetenv ("LIBRARY_PATH");
+ char *omp_requires_file;
+ if (save_temps)
+ omp_requires_file = concat (dumppfx, ".mkoffload.omp_requires", NULL);
+ else
+ omp_requires_file = make_temp_file (".mkoffload.omp_requires");
+
/* Run the compiler pass. */
+ xputenv (concat ("GCC_OFFLOAD_OMP_REQUIRES_FILE=", omp_requires_file, NULL));
fork_execute (cc_argv[0], CONST_CAST (char **, cc_argv), true, ".gcc_args");
obstack_free (&cc_argv_obstack, NULL);
+ unsetenv("GCC_OFFLOAD_OMP_REQUIRES_FILE");
+
+ in = fopen (omp_requires_file, "rb");
+ if (!in)
+ fatal_error (input_location, "cannot open omp_requires file %qs",
+ omp_requires_file);
+ uint32_t omp_requires;
+ if (fread (&omp_requires, sizeof (omp_requires), 1, in) != 1)
+ fatal_error (input_location, "cannot read omp_requires file %qs",
+ omp_requires_file);
+ fclose (in);
in = fopen (gcn_s1_name, "r");
if (!in)
@@ -1102,7 +1123,7 @@ main (int argc, char **argv)
if (!in)
fatal_error (input_location, "cannot open intermediate gcn obj file");
- process_obj (in, cfile);
+ process_obj (in, cfile, omp_requires);
fclose (in);
diff --git a/gcc/config/nvptx/mkoffload.cc b/gcc/config/nvptx/mkoffload.cc
index b28c1a32292a62b6b871c9e656091a50a252dd30..d8c81eb05473994cccc1670092640db0a83e9131 100644
--- a/gcc/config/nvptx/mkoffload.cc
+++ b/gcc/config/nvptx/mkoffload.cc
@@ -231,7 +231,7 @@ access_check (const char *name, int mode)
}
static void
-process (FILE *in, FILE *out)
+process (FILE *in, FILE *out, uint32_t omp_requires)
{
size_t len = 0;
const char *input = read_file (in, &len);
@@ -240,6 +240,8 @@ process (FILE *in, FILE *out)
unsigned obj_count = 0;
unsigned ix;
+ fprintf (out, "#include <stdint.h>\n\n");
+
/* Dump out char arrays for each PTX object file. These are
terminated by a NUL. */
for (size_t i = 0; i != len;)
@@ -309,6 +311,7 @@ process (FILE *in, FILE *out)
fprintf (out,
"static const struct nvptx_tdata {\n"
+ " uintptr_t omp_requires_mask;\n"
" const struct ptx_obj *ptx_objs;\n"
" unsigned ptx_num;\n"
" const char *const *var_names;\n"
@@ -316,12 +319,12 @@ process (FILE *in, FILE *out)
" const struct nvptx_fn *fn_names;\n"
" unsigned fn_num;\n"
"} target_data = {\n"
- " ptx_objs, sizeof (ptx_objs) / sizeof (ptx_objs[0]),\n"
+ " %d, ptx_objs, sizeof (ptx_objs) / sizeof (ptx_objs[0]),\n"
" var_mappings,"
" sizeof (var_mappings) / sizeof (var_mappings[0]),\n"
" func_mappings,"
" sizeof (func_mappings) / sizeof (func_mappings[0])\n"
- "};\n\n");
+ "};\n\n", omp_requires);
fprintf (out, "#ifdef __cplusplus\n"
"extern \"C\" {\n"
@@ -583,19 +586,37 @@ main (int argc, char **argv)
unsetenv ("COMPILER_PATH");
unsetenv ("LIBRARY_PATH");
+ char *omp_requires_file;
+ if (save_temps)
+ omp_requires_file = concat (dumppfx, ".mkoffload.omp_requires", NULL);
+ else
+ omp_requires_file = make_temp_file (".mkoffload.omp_requires");
+
+ xputenv (concat ("GCC_OFFLOAD_OMP_REQUIRES_FILE=", omp_requires_file, NULL));
fork_execute (new_argv[0], CONST_CAST (char **, new_argv), true,
".gcc_args");
obstack_free (&argv_obstack, NULL);
+ unsetenv("GCC_OFFLOAD_OMP_REQUIRES_FILE");
xputenv (concat ("GCC_EXEC_PREFIX=", execpath, NULL));
xputenv (concat ("COMPILER_PATH=", cpath, NULL));
xputenv (concat ("LIBRARY_PATH=", lpath, NULL));
+ in = fopen (omp_requires_file, "rb");
+ if (!in)
+ fatal_error (input_location, "cannot open omp_requires file %qs",
+ omp_requires_file);
+ uint32_t omp_requires;
+ if (fread (&omp_requires, sizeof (omp_requires), 1, in) != 1)
+ fatal_error (input_location, "cannot read omp_requires file %qs",
+ omp_requires_file);
+ fclose (in);
+
in = fopen (ptx_name, "r");
if (!in)
fatal_error (input_location, "cannot open intermediate ptx file");
- process (in, out);
+ process (in, out, omp_requires);
fclose (in);
}
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 6b3763bca5369aa86d541f96f6c3c80add3b2af2..df657a3fb2b1e7e8d0410c3070306e793a0d4bb2 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -44329,6 +44329,10 @@ cp_parser_omp_teams (cp_parser *parser, cp_token *pragma_tok,
static tree
cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
{
+ if (flag_openmp)
+ omp_requires_mask
+ = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
tree clauses
= cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
"#pragma omp target data", pragma_tok);
@@ -44432,6 +44436,10 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok,
return true;
}
+ if (flag_openmp)
+ omp_requires_mask
+ = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
tree clauses
= cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
"#pragma omp target enter data", pragma_tok);
@@ -44531,6 +44539,10 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
return true;
}
+ if (flag_openmp)
+ omp_requires_mask
+ = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
tree clauses
= cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
"#pragma omp target exit data", pragma_tok);
@@ -44625,6 +44637,10 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok,
return true;
}
+ if (flag_openmp)
+ omp_requires_mask
+ = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
tree stmt = make_node (OMP_TARGET_UPDATE);
TREE_TYPE (stmt) = void_type_node;
OMP_TARGET_UPDATE_CLAUSES (stmt) = clauses;
@@ -46919,9 +46935,6 @@ cp_parser_omp_requires (cp_parser *parser, cp_token *pragma_tok)
cp_parser_skip_to_pragma_eol (parser, pragma_tok);
return false;
}
- if (p && this_req != OMP_REQUIRES_DYNAMIC_ALLOCATORS)
- sorry_at (cloc, "%qs clause on %<requires%> directive not "
- "supported yet", p);
if (p)
cp_lexer_consume_token (parser->lexer);
if (this_req)
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 93e40f25f8210212eea2a282839c0f8a3f07addb..51b429a597c07d0199e605a3f46915b8e456d61f 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -5488,10 +5488,6 @@ gfc_match_omp_requires (void)
else
goto error;
- if (requires_clause & ~(OMP_REQ_ATOMIC_MEM_ORDER_MASK
- | OMP_REQ_DYNAMIC_ALLOCATORS))
- gfc_error_now ("Sorry, %qs clause at %L on REQUIRES directive is not "
- "yet supported", clause, &old_loc);
if (!gfc_omp_requires_add_clause (requires_clause, clause, &old_loc, NULL))
goto error;
requires_clauses |= requires_clause;
diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc
index 7356d1b5a3ab206cd4c70ab43776a7e2d1a66148..0b4c596996c6e1406134192beaf7270495e8f42e 100644
--- a/gcc/fortran/parse.cc
+++ b/gcc/fortran/parse.cc
@@ -1168,7 +1168,8 @@ decode_omp_directive (void)
}
switch (ret)
{
- case ST_OMP_DECLARE_TARGET:
+ /* Set omp_target_seen; exclude ST_OMP_DECLARE_TARGET.
+ FIXME: Get clarification, cf. OpenMP Spec Issue #3240. */
case ST_OMP_TARGET:
case ST_OMP_TARGET_DATA:
case ST_OMP_TARGET_ENTER_DATA:
@@ -6879,11 +6880,14 @@ done:
/* Fixup for external procedures and resolve 'omp requires'. */
int omp_requires;
+ bool omp_target_seen;
omp_requires = 0;
+ omp_target_seen = false;
for (gfc_current_ns = gfc_global_ns_list; gfc_current_ns;
gfc_current_ns = gfc_current_ns->sibling)
{
omp_requires |= gfc_current_ns->omp_requires;
+ omp_target_seen |= gfc_current_ns->omp_target_seen;
gfc_check_externals (gfc_current_ns);
}
for (gfc_current_ns = gfc_global_ns_list; gfc_current_ns;
@@ -6908,6 +6912,22 @@ done:
break;
}
+ if (omp_target_seen)
+ omp_requires_mask = (enum omp_requires) (omp_requires_mask
+ | OMP_REQUIRES_TARGET_USED);
+ if (omp_requires & OMP_REQ_REVERSE_OFFLOAD)
+ omp_requires_mask = (enum omp_requires) (omp_requires_mask
+ | OMP_REQUIRES_REVERSE_OFFLOAD);
+ if (omp_requires & OMP_REQ_UNIFIED_ADDRESS)
+ omp_requires_mask = (enum omp_requires) (omp_requires_mask
+ | OMP_REQUIRES_UNIFIED_ADDRESS);
+ if (omp_requires & OMP_REQ_UNIFIED_SHARED_MEMORY)
+ omp_requires_mask
+ = (enum omp_requires) (omp_requires_mask
+ | OMP_REQUIRES_UNIFIED_SHARED_MEMORY);
+ if (omp_requires & OMP_REQ_DYNAMIC_ALLOCATORS)
+ omp_requires_mask = (enum omp_requires) (omp_requires_mask
+ | OMP_REQUIRES_DYNAMIC_ALLOCATORS);
/* Do the parse tree dump. */
gfc_current_ns = flag_dump_fortran_original ? gfc_global_ns_list : NULL;
diff --git a/gcc/lto-cgraph.cc b/gcc/lto-cgraph.cc
index 237743ef0ba5b083b5973a987d7eb530f97b7cf3..48629651e3174df743a0d54967591bd6e2e09d84 100644
--- a/gcc/lto-cgraph.cc
+++ b/gcc/lto-cgraph.cc
@@ -37,6 +37,7 @@ along with GCC; see the file COPYING3. If not see
#include "pass_manager.h"
#include "ipa-utils.h"
#include "omp-offload.h"
+#include "omp-general.h"
#include "stringpool.h"
#include "attribs.h"
#include "alloc-pool.h"
@@ -1068,7 +1069,10 @@ read_string (class lto_input_block *ib)
void
output_offload_tables (void)
{
- if (vec_safe_is_empty (offload_funcs) && vec_safe_is_empty (offload_vars))
+ bool output_requires = (flag_openmp
+ && (omp_requires_mask & OMP_REQUIRES_TARGET_USED) != 0);
+ if (vec_safe_is_empty (offload_funcs) && vec_safe_is_empty (offload_vars)
+ && !output_requires)
return;
struct lto_simple_output_block *ob
@@ -1098,6 +1102,19 @@ output_offload_tables (void)
(*offload_vars)[i]);
}
+ if (output_requires)
+ {
+ HOST_WIDE_INT val = ((HOST_WIDE_INT) omp_requires_mask
+ & (OMP_REQUIRES_UNIFIED_ADDRESS
+ | OMP_REQUIRES_UNIFIED_SHARED_MEMORY
+ | OMP_REQUIRES_REVERSE_OFFLOAD
+ | OMP_REQUIRES_TARGET_USED));
+ /* (Mis)use LTO_symtab_edge for this variable. */
+ streamer_write_enum (ob->main_stream, LTO_symtab_tags,
+ LTO_symtab_last_tag, LTO_symtab_edge);
+ streamer_write_hwi_stream (ob->main_stream, val);
+ }
+
streamer_write_uhwi_stream (ob->main_stream, 0);
lto_destroy_simple_output_block (ob);
@@ -1764,6 +1781,20 @@ input_symtab (void)
}
}
+static void
+omp_requires_to_name (char *buf, size_t size, HOST_WIDE_INT requires_mask)
+{
+ char *end = buf + size, *p = buf;
+ if (requires_mask & GOMP_REQUIRES_UNIFIED_ADDRESS)
+ p += snprintf (p, end - p, "unified_address");
+ if (requires_mask & GOMP_REQUIRES_UNIFIED_SHARED_MEMORY)
+ p += snprintf (p, end - p, "%sunified_shared_memory",
+ (p == buf ? "" : ", "));
+ if (requires_mask & GOMP_REQUIRES_REVERSE_OFFLOAD)
+ p += snprintf (p, end - p, "%sreverse_offload",
+ (p == buf ? "" : ", "));
+}
+
/* Input function/variable tables that will allow libgomp to look up offload
target code, and store them into OFFLOAD_FUNCS and OFFLOAD_VARS. */
@@ -1773,6 +1804,10 @@ input_offload_tables (bool do_force_output)
struct lto_file_decl_data **file_data_vec = lto_get_file_decl_data ();
struct lto_file_decl_data *file_data;
unsigned int j = 0;
+ const char *requires_fn = NULL;
+ tree requires_decl = NULL_TREE;
+
+ omp_requires_mask = (omp_requires) 0;
while ((file_data = file_data_vec[j++]))
{
@@ -1784,6 +1819,7 @@ input_offload_tables (bool do_force_output)
if (!ib)
continue;
+ tree tmp_decl = NULL_TREE;
enum LTO_symtab_tags tag
= streamer_read_enum (ib, LTO_symtab_tags, LTO_symtab_last_tag);
while (tag)
@@ -1799,6 +1835,7 @@ input_offload_tables (bool do_force_output)
LTO mode. */
if (do_force_output)
cgraph_node::get (fn_decl)->mark_force_output ();
+ tmp_decl = fn_decl;
}
else if (tag == LTO_symtab_variable)
{
@@ -1810,6 +1847,72 @@ input_offload_tables (bool do_force_output)
may be no refs to var_decl in offload LTO mode. */
if (do_force_output)
varpool_node::get (var_decl)->force_output = 1;
+ tmp_decl = var_decl;
+ }
+ else if (tag == LTO_symtab_edge)
+ {
+ static bool error_emitted = false;
+ HOST_WIDE_INT val = streamer_read_hwi (ib);
+
+ if (omp_requires_mask == 0)
+ {
+ omp_requires_mask = (omp_requires) val;
+ requires_decl = tmp_decl;
+ requires_fn = file_data->file_name;
+ }
+ else if (omp_requires_mask != val && !error_emitted)
+ {
+ const char *fn1 = requires_fn;
+ if (requires_decl != NULL_TREE)
+ {
+ while (DECL_CONTEXT (requires_decl) != NULL_TREE
+ && TREE_CODE (requires_decl) != TRANSLATION_UNIT_DECL)
+ requires_decl = DECL_CONTEXT (requires_decl);
+ if (requires_decl != NULL_TREE)
+ fn1 = IDENTIFIER_POINTER (DECL_NAME (requires_decl));
+ }
+
+ const char *fn2 = file_data->file_name;
+ if (tmp_decl != NULL_TREE)
+ {
+ while (DECL_CONTEXT (tmp_decl) != NULL_TREE
+ && TREE_CODE (tmp_decl) != TRANSLATION_UNIT_DECL)
+ tmp_decl = DECL_CONTEXT (tmp_decl);
+ if (tmp_decl != NULL_TREE)
+ fn2 = IDENTIFIER_POINTER (DECL_NAME (requires_decl));
+ }
+
+ char buf1[sizeof ("unified_address, unified_shared_memory, "
+ "reverse_offload")];
+ char buf2[sizeof ("unified_address, unified_shared_memory, "
+ "reverse_offload")];
+ omp_requires_to_name (buf2, sizeof (buf2),
+ val != OMP_REQUIRES_TARGET_USED
+ ? val
+ : (HOST_WIDE_INT) omp_requires_mask);
+ if (val != OMP_REQUIRES_TARGET_USED
+ && omp_requires_mask != OMP_REQUIRES_TARGET_USED)
+ {
+ omp_requires_to_name (buf1, sizeof (buf1),
+ omp_requires_mask);
+ error ("OpenMP %<requires%> directive with non-identical "
+ "clauses in multiple compilation units: %qs vs. "
+ "%qs", buf1, buf2);
+ inform (UNKNOWN_LOCATION, "%qs has %qs", fn1, buf1);
+ inform (UNKNOWN_LOCATION, "%qs has %qs", fn2, buf2);
+ }
+ else
+ {
+ error ("OpenMP %<requires%> directive with %qs specified "
+ "only in some compilation units", buf2);
+ inform (UNKNOWN_LOCATION, "%qs has %qs",
+ val != OMP_REQUIRES_TARGET_USED ? fn2 : fn1,
+ buf2);
+ inform (UNKNOWN_LOCATION, "but %qs has not",
+ val != OMP_REQUIRES_TARGET_USED ? fn1 : fn2);
+ }
+ error_emitted = true;
+ }
}
else
fatal_error (input_location,
@@ -1821,6 +1924,18 @@ input_offload_tables (bool do_force_output)
lto_destroy_simple_input_block (file_data, LTO_section_offload_table,
ib, data, len);
}
+#ifdef ACCEL_COMPILER
+ char *omp_requires_file = getenv ("GCC_OFFLOAD_OMP_REQUIRES_FILE");
+ if (omp_requires_file == NULL || omp_requires_file[0] == '\0')
+ fatal_error (input_location, "GCC_OFFLOAD_OMP_REQUIRES_FILE unset");
+ FILE *f = fopen (omp_requires_file, "wb");
+ if (!f)
+ fatal_error (input_location, "Cannot open omp_requires file %qs",
+ omp_requires_file);
+ uint32_t req_mask = omp_requires_mask;
+ fwrite (&req_mask, sizeof (req_mask), 1, f);
+ fclose (f);
+#endif
}
/* True when we need optimization summary for NODE. */
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index b9d5529f212773379e7de430ff0dfb60d6398416..d73c165f0298a6098829ddd2d8a8ce055b0681dd 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -12701,6 +12701,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gcc_unreachable ();
}
+ /* Ensure that requires map is written via output_offload_tables, even if only
+ 'target (enter/exit) data' is used in the translation unit. */
+ if (ENABLE_OFFLOADING && (omp_requires_mask & OMP_REQUIRES_TARGET_USED))
+ g->have_offload = true;
+
clauses = gimple_omp_target_clauses (stmt);
gimple_seq dep_ilist = NULL;
diff --git a/gcc/testsuite/c-c++-common/gomp/requires-4.c b/gcc/testsuite/c-c++-common/gomp/requires-4.c
index 88ba7746cf81698c3aeb3f9ed877448fd5ca43e0..8f45d83ea6e30c2d564ec14d6fd5f463575e1a51 100644
--- a/gcc/testsuite/c-c++-common/gomp/requires-4.c
+++ b/gcc/testsuite/c-c++-common/gomp/requires-4.c
@@ -9,5 +9,3 @@ foo (void)
#pragma omp requires unified_shared_memory /* { dg-error "'unified_shared_memory' clause used lexically after first target construct or offloading API" } */
#pragma omp requires unified_address /* { dg-error "'unified_address' clause used lexically after first target construct or offloading API" } */
#pragma omp requires reverse_offload /* { dg-error "'reverse_offload' clause used lexically after first target construct or offloading API" } */
-
-/* { dg-prune-output "not supported yet" } */
diff --git a/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c b/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c
index 9a3fa5230f8b956e6a5bba676d8e4951b2f3c0d9..3452156f9484ea0e99d12aa2fc7e24af3a1d095f 100644
--- a/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c
@@ -43,7 +43,7 @@ tg_fn (int *x, int *y)
x2 = x2 + 2 + called_in_target1 ();
y2 = y2 + 7;
- #pragma omp target device(ancestor : 1) map(tofrom: x2)
+ #pragma omp target device(ancestor : 1) map(tofrom: x2) /* { dg-message "sorry, unimplemented: 'ancestor' not yet supported" } */
check_offload(&x2, &y2);
if (x2 != 2+2+3+42 || y2 != 3 + 7)
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c
index cf05c505004b3352d4562c6471376859c6957c34..b16e701bd5a175b5d647143e15fe35b9a32ad89b 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c
@@ -1,13 +1,11 @@
/* { dg-do compile } */
-#pragma omp requires reverse_offload /* { dg-message "sorry, unimplemented: 'reverse_offload' clause on 'requires' directive not supported yet" } */
+#pragma omp requires reverse_offload
void
foo (int n)
{
- /* The following test is marked with 'xfail' because a previous 'sorry' from
- 'reverse_offload' suppresses the 'sorry' for 'ancestor'. */
- #pragma omp target device (ancestor: 1) /* { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } } */
+ #pragma omp target device (ancestor: 1)
;
@@ -19,9 +17,9 @@ foo (int n)
#pragma omp target device (ancestor : 42) /* { dg-error "the 'device' clause expression must evaluate to '1'" } */
;
- #pragma omp target device (ancestor : n) /* { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } } */
+ #pragma omp target device (ancestor : n)
;
- #pragma omp target device (ancestor : n + 1) /* { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } } */
+ #pragma omp target device (ancestor : n + 1)
;
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-3.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-3.c
index ea6e5a0cf6c5e1d2c67714b648234492001d3eb2..d16590107d2fe461f2b220bcf522340e690ddafc 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-3.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-3.c
@@ -11,7 +11,7 @@ int bar (void);
/* { dg-do compile } */
-#pragma omp requires reverse_offload /* { dg-message "sorry, unimplemented: 'reverse_offload' clause on 'requires' directive not supported yet" } */
+#pragma omp requires reverse_offload
void
foo (void)
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-4.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-4.c
index b4b5620bbc0f90ab93b33a1b9b10911492b2d541..241234f8daf6d25d36c396ad3c1ad4a28acc4e2a 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-4.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-4.c
@@ -4,12 +4,12 @@
/* Test to ensure that device-modifier 'ancestor' is parsed correctly in
device clauses. */
-#pragma omp requires reverse_offload /* { dg-message "sorry, unimplemented: 'reverse_offload' clause on 'requires' directive not supported yet" } */
+#pragma omp requires reverse_offload
void
foo (void)
{
- #pragma omp target device (ancestor: 1) /* { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } } */
+ #pragma omp target device (ancestor: 1) /* { dg-message "sorry, unimplemented: 'ancestor' not yet supported" } */
;
}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c
index b6ff84bcdab9d5f1f9e7161ce3a86f0fba505471..b1520ff0636e415dadd9957b4dae1d818bd82e38 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c
@@ -1,4 +1,4 @@
-#pragma omp requires reverse_offload /* { dg-message "sorry, unimplemented: 'reverse_offload' clause on 'requires' directive not supported yet" } */
+#pragma omp requires reverse_offload
void
foo ()
diff --git a/gcc/testsuite/gfortran.dg/gomp/requires-8.f90 b/gcc/testsuite/gfortran.dg/gomp/requires-8.f90
index e84d609ad298cf79ce66e0284d66f61436bb7290..583c5a56b32e07f7c1f50aea589ce4012a73cd2b 100644
--- a/gcc/testsuite/gfortran.dg/gomp/requires-8.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/requires-8.f90
@@ -1,3 +1,7 @@
+module m0
+ integer :: x
+end module m0
+
module m ! { dg-error "has OpenMP device constructs/routines but does not set !.OMP REQUIRES UNIFIED_SHARED_MEMORY but other program units do" }
!$omp requires reverse_offload
contains
@@ -13,10 +17,14 @@ contains
end subroutine foo
end module m
-subroutine bar ! { dg-error "has OpenMP device constructs/routines but does not set !.OMP REQUIRES REVERSE_OFFLOAD but other program units do" }
+subroutine bar
!use m
- !$omp requires unified_shared_memory
+ !$omp requires unified_shared_memory ! Possibly OK - needs OpenMP Lang Spec clarification (-> #3240)
!$omp declare target
end subroutine bar
-! { dg-prune-output "not yet supported" }
+subroutine foobar ! { dg-error "has OpenMP device constructs/routines but does not set !.OMP REQUIRES REVERSE_OFFLOAD but other program units do" }
+ use m0
+ !$omp requires unified_shared_memory
+ !$omp target enter data map(to:x)
+end subroutine foobar
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2.f90
index 117a1d000a5748fb6784168132ea6f9458bd1d4d..230c690d84c5f94d05e58e1ef9879a6a63094644 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2.f90
@@ -4,19 +4,16 @@ implicit none
integer :: a, b, c
-!$omp requires reverse_offload ! { dg-error "Sorry, 'reverse_offload' clause at \\(1\\) on REQUIRES directive is not yet supported" }
+!$omp requires reverse_offload
-! The following test case is marked with 'xfail' because a previous 'sorry' from
-! 'reverse_offload' suppresses the 'sorry' for 'ancestor'.
-
-!$omp target device (ancestor: 1) ! { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } }
+!$omp target device (ancestor: 1)
!$omp end target
-!$omp target device (ancestor : a) ! { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } }
+!$omp target device (ancestor : a)
!$omp end target
-!$omp target device (ancestor : a + 1) ! { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } }
+!$omp target device (ancestor : a + 1)
!$omp end target
@@ -32,61 +29,4 @@ integer :: a, b, c
!$omp target device (42)
!$omp end target
-
-! Ensure that no OpenMP constructs appear inside target regions with 'ancestor'.
-! The following test case is marked with 'xfail' because a previous 'sorry' from
-! 'reverse_offload' suppresses the 'sorry' for 'ancestor'.
-
-!$omp target device (ancestor: 1)
- !$omp teams ! { dg-error "" "OpenMP constructs are not allowed in target region with 'ancestor'" { xfail *-*-* } }
- !$omp end teams
-!$omp end target
-
-!$omp target device (device_num: 1)
- !$omp teams
- !$omp end teams
-!$omp end target
-
-!$omp target device (1)
- !$omp teams
- !$omp end teams
-!$omp end target
-
-
-! Ensure that with 'ancestor' only the 'device', 'firstprivate', 'private',
-! 'defaultmap', and 'map' clauses appear on the construct.
-! The following test case is marked with 'xfail' because a previous 'sorry' from
-! 'reverse_offload' suppresses the 'sorry' for 'ancestor'.
-
-!$omp target nowait device (ancestor: 1) ! { dg-error "" "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" { xfail *-*-* } }
-!$omp end target
-
-!$omp target device (ancestor: 1) nowait ! { dg-error "" "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" { xfail *-*-* } }
-!$omp end target
-
-!$omp target nowait device (device_num: 1)
-!$omp end target
-
-!$omp target nowait device (1)
-!$omp end target
-
-!$omp target device (ancestor: 1) firstprivate (a) private (b) defaultmap (none) map (c)
-!$omp end target
-
-
-! Ensure that 'ancestor' is only used with 'target' constructs (not with
-! 'target data', 'target update' etc.).
-! The following test case is marked with 'xfail' because a previous 'sorry' from
-! 'reverse_offload' suppresses the 'sorry' for 'ancestor'.
-
-!$omp target data map (a) device (ancestor: 1) ! { dg-error "" "'device' clause with 'ancestor' is only allowed on 'target' construct" { xfail *-*-* } }
-!$omp end target data
-
-!$omp target enter data map (to: a) device (ancestor: 1) ! { dg-error "" "'device' clause with 'ancestor' is only allowed on 'target' construct" { xfail *-*-* } }
-!$omp target exit data map (from: a) device (ancestor: 1) ! { dg-error "" "'device' clause with 'ancestor' is only allowed on 'target' construct" { xfail *-*-* } }
-
-!$omp target update to (a) device (ancestor: 1) ! { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" "" { xfail *-*-* } }
-! { dg-error "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" "" { xfail *-*-* } .-1 }
-
-
-end
\ No newline at end of file
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2a.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2a.f90
new file mode 100644
index 0000000000000000000000000000000000000000..feb76fe214419e94f2a0b9f21128f1f6e32a0a16
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2a.f90
@@ -0,0 +1,80 @@
+! { dg-do compile }
+
+implicit none
+
+integer :: a, b, c
+
+!$omp requires reverse_offload
+
+!$omp target device (ancestor: 1)
+!$omp end target
+
+!$omp target device (ancestor : a)
+!$omp end target
+
+!$omp target device (ancestor : a + 1)
+!$omp end target
+
+
+!$omp target device (device_num:42)
+!$omp end target
+
+!$omp target device (42)
+!$omp end target
+
+
+! Ensure that no OpenMP constructs appear inside target regions with 'ancestor'.
+
+!$omp target device (ancestor: 1)
+ !$omp teams ! { dg-error "OpenMP constructs are not allowed in target region with 'ancestor'" }
+ !$omp end teams
+!$omp end target
+
+!$omp target device (device_num: 1)
+ !$omp teams
+ !$omp end teams
+!$omp end target
+
+!$omp target device (1)
+ !$omp teams
+ !$omp end teams
+!$omp end target
+
+
+! Ensure that with 'ancestor' only the 'device', 'firstprivate', 'private',
+! 'defaultmap', and 'map' clauses appear on the construct.
+
+!$omp target nowait device (ancestor: 1) ! { dg-error "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" }
+!$omp end target
+
+!$omp target device (ancestor: 1) nowait ! { dg-error "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" }
+!$omp end target
+
+!$omp target nowait device (device_num: 1)
+!$omp end target
+
+!$omp target nowait device (1)
+!$omp end target
+
+!$omp target device (ancestor: 1) firstprivate (a) private (b) defaultmap (none) map (c)
+!$omp end target
+
+
+! Ensure that 'ancestor' is only used with 'target' constructs (not with
+! 'target data', 'target update' etc.).
+! The following test case is marked with 'xfail' because a previous 'sorry' from
+! 'reverse_offload' suppresses the 'sorry' for 'ancestor'.
+
+!$omp target data map (a) device (ancestor: 1) ! { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" }
+!$omp end target data
+
+!$omp target enter data map (to: a) device (ancestor: 1) ! { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" }
+!$omp target exit data map (from: a) device (ancestor: 1) ! { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" }
+
+!$omp target update to (a) device (ancestor: 1) ! { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" }
+
+!$omp target device (ancestor: 1) if(.false.)
+! { dg-error "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" "" { target *-*-* } .-1 }
+!$omp end target
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-3.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-3.f90
index f1145bde2ece0bee085c97f68de4b64be032b984..e8975e6a08b5047d169e68ac673c7a9e4d7ecc8f 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-3.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-3.f90
@@ -16,10 +16,10 @@ subroutine f1 ()
implicit none
integer :: n
- !$omp requires reverse_offload ! { dg-error "Sorry, 'reverse_offload' clause at \\(1\\) on REQUIRES directive is not yet supported" }
+ !$omp requires reverse_offload
!$omp target device (ancestor : 1)
- n = omp_get_thread_num () ! { dg-error "" "OpenMP runtime API call 'omp_get_thread_num' in a region with 'device\\(ancestor\\)' clause" { xfail *-*-* } }
+ n = omp_get_thread_num () ! { dg-error "OpenMP runtime API call 'omp_get_thread_num' in a region with 'device\\(ancestor\\)' clause" }
!$omp end target
!$omp target device (device_num : 1)
@@ -30,4 +30,4 @@ subroutine f1 ()
n = omp_get_thread_num ()
!$omp end target
-end
\ No newline at end of file
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-4.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-4.f90
index 63872fa51fb5cca2d4c4708c6b63c2086445a199..ab56e2d1d52a654108f2aeae184567593ed9caef 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-4.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-4.f90
@@ -4,11 +4,11 @@
! Test to ensure that device-modifier 'ancestor' is parsed correctly in
! device clauses.
-!$omp requires reverse_offload ! { dg-error "Sorry, 'reverse_offload' clause at \\(1\\) on REQUIRES directive is not yet supported" }
+!$omp requires reverse_offload
-!$omp target device (ancestor : 1) ! { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } }
+!$omp target device (ancestor : 1) ! { dg-message "sorry, unimplemented: 'ancestor' not yet supported" }
!$omp end target
end
-! TODO: dg-final { scan-tree-dump-times "pragma omp target \[^\n\r)]*device\\(ancestor:1\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "pragma omp target \[^\n\r)]*device\\(ancestor:1\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-5.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-5.f90
index 06a11eb5092bd2357c343d920f47dfb5a75c8035..ca8d4b282a0d9842a3be3560d10be815c5c11a42 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-5.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-5.f90
@@ -6,7 +6,7 @@
!
module m
- !$omp requires reverse_offload ! { dg-error "REQUIRES directive is not yet supported" }
+ !$omp requires reverse_offload
contains
subroutine foo()
!$omp target device(ancestor:1)
@@ -17,7 +17,7 @@ contains
block
block
block
- !$omp target device(ancestor:1)
+ !$omp target device(ancestor:1) ! { dg-message "sorry, unimplemented: 'ancestor' not yet supported" }
!$omp end target
end block
end block
@@ -26,7 +26,7 @@ contains
end module m
subroutine foo()
- !$omp requires reverse_offload ! { dg-error "REQUIRES directive is not yet supported" }
+ !$omp requires reverse_offload
block
block
block
@@ -49,7 +49,7 @@ contains
end subroutine foo
program main
- !$omp requires reverse_offload ! { dg-error "REQUIRES directive is not yet supported" }
+ !$omp requires reverse_offload
contains
subroutine foo()
!$omp target device(ancestor:1)
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index e4dd8ef3e1d892edbc80168081e93bbec37a66f1..3e3078f082ed3f1cd9b42bc51f0c61f950c627f7 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -282,7 +282,7 @@ enum gomp_map_kind
/* Versions of libgomp and device-specific plugins. GOMP_VERSION
should be incremented whenever an ABI-incompatible change is introduced
to the plugin interface defined in libgomp/libgomp.h. */
-#define GOMP_VERSION 1
+#define GOMP_VERSION 2
#define GOMP_VERSION_NVIDIA_PTX 1
#define GOMP_VERSION_INTEL_MIC 0
#define GOMP_VERSION_GCN 2
@@ -341,6 +341,13 @@ enum gomp_map_kind
#define GOMP_DEPEND_MUTEXINOUTSET 4
#define GOMP_DEPEND_INOUTSET 5
+/* Flag values for requires-directive features, must match corresponding
+ OMP_REQUIRES_* values in gcc/omp-general.h. */
+#define GOMP_REQUIRES_UNIFIED_ADDRESS 0x10
+#define GOMP_REQUIRES_UNIFIED_SHARED_MEMORY 0x20
+#define GOMP_REQUIRES_REVERSE_OFFLOAD 0x80
+#define GOMP_REQUIRES_TARGET_USED 0x200
+
/* HSA specific data structures. */
/* Identifiers of device-specific target arguments. */
diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h
index 07ab700b80cc62676926f4ae5ce83b1d561c156e..ab3ed638475be3ab544fbd9a84d5c22349d24ac5 100644
--- a/libgomp/libgomp-plugin.h
+++ b/libgomp/libgomp-plugin.h
@@ -125,7 +125,7 @@ extern void GOMP_PLUGIN_fatal (const char *, ...)
extern const char *GOMP_OFFLOAD_get_name (void);
extern unsigned int GOMP_OFFLOAD_get_caps (void);
extern int GOMP_OFFLOAD_get_type (void);
-extern int GOMP_OFFLOAD_get_num_devices (void);
+extern int GOMP_OFFLOAD_get_num_devices (unsigned int);
extern bool GOMP_OFFLOAD_init_device (int);
extern bool GOMP_OFFLOAD_fini_device (int);
extern unsigned GOMP_OFFLOAD_version (void);
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index a75cd244a8397ab398111685e62f2210cdaf9857..39426ff7fbfbc345e16e7e70f916f8390ca5048d 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -189,8 +189,8 @@ The OpenMP 4.5 specification is fully supported.
env variable @tab Y @tab
@item Nested-parallel changes to @emph{max-active-levels-var} ICV @tab Y @tab
@item @code{requires} directive @tab P
- @tab Only fulfillable requirement are @code{atomic_default_mem_order}
- and @code{dynamic_allocators}
+ @tab complete but no non-host devices provides @code{unified_address},
+ @code{unified_shared_memory} or @code{reverse_offload}
@item @code{teams} construct outside an enclosing target region @tab Y @tab
@item Non-rectangular loop nests @tab Y @tab
@item @code{!=} as relational-op in canonical loop form for C/C++ @tab Y @tab
@@ -344,6 +344,8 @@ The OpenMP 4.5 specification is fully supported.
@item @code{unconstrained} and @code{reproducible} modifiers on @code{order}
clause @tab Y @tab
@item Support @code{begin/end declare target} syntax in C/C++ @tab N @tab
+@item Pointer predetermined firstprivate getting initialized
+to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab
@end multitable
@@ -361,7 +363,7 @@ The OpenMP 4.5 specification is fully supported.
@item Clauses on @code{end} directive can be on directive @tab N @tab
@item Deprecation of no-argument @code{destroy} clause on @code{depobj}
@tab N @tab
-@item @code{linear} clause syntax changes and @code{step} modifier @tab N @tab
+@item @code{linear} clause syntax changes and @code{step} modifier @tab P @tab only C/C++
@item Deprecation of minus operator for reductions @tab N @tab
@item Deprecation of separating @code{map} modifiers without comma @tab N @tab
@item @code{declare mapper} with iterator and @code{present} modifiers
diff --git a/libgomp/oacc-host.c b/libgomp/oacc-host.c
index 5bb889926d3303d7dfd9c052ba71d255df10ea19..eb11b9cf16a90780c2796a82614d58c05640245a 100644
--- a/libgomp/oacc-host.c
+++ b/libgomp/oacc-host.c
@@ -54,7 +54,7 @@ host_get_type (void)
}
static int
-host_get_num_devices (void)
+host_get_num_devices (unsigned int omp_requires_mask __attribute__((unused)))
{
return 1;
}
@@ -229,7 +229,7 @@ host_openacc_get_property (int n, enum goacc_property prop)
{
union goacc_property_value nullval = { .val = 0 };
- if (n >= host_get_num_devices ())
+ if (n >= host_get_num_devices (0))
return nullval;
switch (prop)
diff --git a/libgomp/oacc-init.c b/libgomp/oacc-init.c
index 1565aa0f290bbd4afa6f49fe483d66f3cc3eaca0..42c3e74e6ba3f7fb4e44696584f04b921ecd1443 100644
--- a/libgomp/oacc-init.c
+++ b/libgomp/oacc-init.c
@@ -148,7 +148,7 @@ resolve_device (acc_device_t d, bool fail_is_error)
if (dispatchers[d]
&& !strcasecmp (goacc_device_type,
get_openacc_name (dispatchers[d]->name))
- && dispatchers[d]->get_num_devices_func () > 0)
+ && dispatchers[d]->get_num_devices_func (0) > 0)
goto found;
if (fail_is_error)
@@ -169,7 +169,7 @@ resolve_device (acc_device_t d, bool fail_is_error)
case acc_device_not_host:
/* Find the first available device after acc_device_not_host. */
while (known_device_type_p (++d))
- if (dispatchers[d] && dispatchers[d]->get_num_devices_func () > 0)
+ if (dispatchers[d] && dispatchers[d]->get_num_devices_func (0) > 0)
goto found;
if (d_arg == acc_device_default)
{
@@ -302,7 +302,7 @@ acc_init_1 (acc_device_t d, acc_construct_t parent_construct, int implicit)
base_dev = resolve_device (d, true);
- ndevs = base_dev->get_num_devices_func ();
+ ndevs = base_dev->get_num_devices_func (0);
if (ndevs <= 0 || goacc_device_num >= ndevs)
acc_dev_num_out_of_range (d, goacc_device_num, ndevs);
@@ -351,7 +351,7 @@ acc_shutdown_1 (acc_device_t d)
/* Get the base device for this device type. */
base_dev = resolve_device (d, true);
- ndevs = base_dev->get_num_devices_func ();
+ ndevs = base_dev->get_num_devices_func (0);
/* Unload all the devices of this type that have been opened. */
for (i = 0; i < ndevs; i++)
@@ -520,7 +520,7 @@ goacc_attach_host_thread_to_device (int ord)
base_dev = cached_base_dev;
}
- num_devices = base_dev->get_num_devices_func ();
+ num_devices = base_dev->get_num_devices_func (0);
if (num_devices <= 0 || ord >= num_devices)
acc_dev_num_out_of_range (acc_device_type (base_dev->type), ord,
num_devices);
@@ -599,7 +599,7 @@ acc_get_num_devices (acc_device_t d)
if (!acc_dev)
return 0;
- n = acc_dev->get_num_devices_func ();
+ n = acc_dev->get_num_devices_func (0);
if (n < 0)
n = 0;
@@ -779,7 +779,7 @@ acc_set_device_num (int ord, acc_device_t d)
cached_base_dev = base_dev = resolve_device (d, true);
- num_devices = base_dev->get_num_devices_func ();
+ num_devices = base_dev->get_num_devices_func (0);
if (num_devices <= 0 || ord >= num_devices)
acc_dev_num_out_of_range (d, ord, num_devices);
@@ -814,7 +814,7 @@ get_property_any (int ord, acc_device_t d, acc_device_property_t prop)
struct gomp_device_descr *dev = resolve_device (d, true);
- int num_devices = dev->get_num_devices_func ();
+ int num_devices = dev->get_num_devices_func (0);
if (num_devices <= 0 || ord >= num_devices)
acc_dev_num_out_of_range (d, ord, num_devices);
diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c
index 1c0436842da5602d17a801b353804e1b4aee3505..ea327bf2ca09c16105b54e390ab9a57ea59774f9 100644
--- a/libgomp/plugin/plugin-gcn.c
+++ b/libgomp/plugin/plugin-gcn.c
@@ -3221,10 +3221,14 @@ GOMP_OFFLOAD_version (void)
/* Return the number of GCN devices on the system. */
int
-GOMP_OFFLOAD_get_num_devices (void)
+GOMP_OFFLOAD_get_num_devices (unsigned int omp_requires_mask)
{
if (!init_hsa_context ())
return 0;
+ /* Return -1 if no omp_requires_mask cannot be fulfilled but
+ devices were present. */
+ if (hsa_context.agent_count > 0 && omp_requires_mask != 0)
+ return -1;
return hsa_context.agent_count;
}
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 387bcbbc52afb2ae22fb91d53358951a89d9b338..bc63e274cdfa7639c8663de9e1d4737fce2a3a4e 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -1175,9 +1175,14 @@ GOMP_OFFLOAD_get_type (void)
}
int
-GOMP_OFFLOAD_get_num_devices (void)
+GOMP_OFFLOAD_get_num_devices (unsigned int omp_requires_mask)
{
- return nvptx_get_num_devices ();
+ int num_devices = nvptx_get_num_devices ();
+ /* Return -1 if no omp_requires_mask cannot be fulfilled but
+ devices were present. */
+ if (num_devices > 0 && omp_requires_mask != 0)
+ return -1;
+ return num_devices;
}
bool
diff --git a/libgomp/target.c b/libgomp/target.c
index c0844f2265a6b3b4c2887d66a0cdf654d55c8876..4dac81862d7f73fece6b05850e60776760e41cd3 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -36,6 +36,7 @@
# include <inttypes.h> /* For PRIu64. */
#endif
#include <string.h>
+#include <stdio.h> /* For snprintf. */
#include <assert.h>
#include <errno.h>
@@ -98,6 +99,9 @@ static int num_devices;
/* Number of GOMP_OFFLOAD_CAP_OPENMP_400 devices. */
static int num_devices_openmp;
+/* OpenMP requires mask. */
+static int omp_requires_mask;
+
/* Similar to gomp_realloc, but release register_lock before gomp_fatal. */
static void *
@@ -2314,6 +2318,20 @@ gomp_unload_image_from_device (struct gomp_device_descr *devicep,
}
}
+static void
+gomp_requires_to_name (char *buf, size_t size, int requires_mask)
+{
+ char *end = buf + size, *p = buf;
+ if (requires_mask & GOMP_REQUIRES_UNIFIED_ADDRESS)
+ p += snprintf (p, end - p, "unified_address");
+ if (requires_mask & GOMP_REQUIRES_UNIFIED_SHARED_MEMORY)
+ p += snprintf (p, end - p, "%sunified_shared_memory",
+ (p == buf ? "" : ", "));
+ if (requires_mask & GOMP_REQUIRES_REVERSE_OFFLOAD)
+ p += snprintf (p, end - p, "%sreverse_offload",
+ (p == buf ? "" : ", "));
+}
+
/* This function should be called from every offload image while loading.
It gets the descriptor of the host func and var tables HOST_TABLE, TYPE of
the target, and TARGET_DATA needed by target plugin. */
@@ -2323,13 +2341,43 @@ GOMP_offload_register_ver (unsigned version, const void *host_table,
int target_type, const void *target_data)
{
int i;
+ int omp_req = 0;
if (GOMP_VERSION_LIB (version) > GOMP_VERSION)
gomp_fatal ("Library too old for offload (version %u < %u)",
GOMP_VERSION, GOMP_VERSION_LIB (version));
-
+
+ if (GOMP_VERSION_LIB (version) > 1)
+ {
+ omp_req = (int) (size_t) ((void **) target_data)[0];
+ target_data = &((void **) target_data)[1];
+ }
+
gomp_mutex_lock (®ister_lock);
+ if (omp_req && omp_requires_mask && omp_requires_mask != omp_req)
+ {
+ char buf1[sizeof ("unified_address, unified_shared_memory, "
+ "reverse_offload")];
+ char buf2[sizeof ("unified_address, unified_shared_memory, "
+ "reverse_offload")];
+ gomp_requires_to_name (buf2, sizeof (buf2),
+ omp_req != GOMP_REQUIRES_TARGET_USED
+ ? omp_req : omp_requires_mask);
+ if (omp_req != GOMP_REQUIRES_TARGET_USED
+ && omp_requires_mask != GOMP_REQUIRES_TARGET_USED)
+ {
+ gomp_requires_to_name (buf1, sizeof (buf1), omp_requires_mask);
+ gomp_fatal ("OpenMP 'requires' directive with non-identical clauses "
+ "in multiple compilation units: '%s' vs. '%s'",
+ buf1, buf2);
+ }
+ else
+ gomp_fatal ("OpenMP 'requires' directive with '%s' specified only in "
+ "some compilation units", buf2);
+ }
+ omp_requires_mask = omp_req;
+
/* Load image to all initialized devices. */
for (i = 0; i < num_devices; i++)
{
@@ -4125,8 +4173,30 @@ gomp_target_init (void)
if (gomp_load_plugin_for_device (¤t_device, plugin_name))
{
- new_num_devs = current_device.get_num_devices_func ();
- if (new_num_devs >= 1)
+ int omp_req = omp_requires_mask & ~GOMP_REQUIRES_TARGET_USED;
+ new_num_devs = current_device.get_num_devices_func (omp_req);
+ if (gomp_debug_var > 0 && new_num_devs < 0)
+ {
+ bool found = false;
+ int type = current_device.get_type_func ();
+ for (int img = 0; img < num_offload_images; img++)
+ if (type == offload_images[img].type)
+ found = true;
+ if (found)
+ {
+ char buf[sizeof ("unified_address, unified_shared_memory, "
+ "reverse_offload")];
+ gomp_requires_to_name (buf, sizeof (buf), omp_req);
+ char *name = (char *) malloc (cur_len + 1);
+ memcpy (name, cur, cur_len);
+ name[cur_len] = '\0';
+ gomp_debug (1,
+ "%s devices present but 'omp requires %s' "
+ "cannot be fulfilled", name, buf);
+ free (name);
+ }
+ }
+ else if (new_num_devs >= 1)
{
/* Augment DEVICES and NUM_DEVICES. */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-1-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-1-aux.c
new file mode 100644
index 0000000000000000000000000000000000000000..bdca662e42f9043224d02d1639e18b3428116b8e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-1-aux.c
@@ -0,0 +1,11 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires unified_address
+
+int x;
+
+void foo (void)
+{
+ #pragma omp target
+ x = 1;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-1.c b/libgomp/testsuite/libgomp.c-c++-common/requires-1.c
new file mode 100644
index 0000000000000000000000000000000000000000..fedf977976954d39ef6df93c9dbe0e4e9ea31de9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-1.c
@@ -0,0 +1,24 @@
+/* { dg-do link { target { offload_target_nvptx || offload_target_amdgcn } } } */
+/* { dg-additional-sources requires-1-aux.c } */
+
+/* Check diagnostic by device-compiler's lto1.
+ Other file uses: 'requires unified_address'. */
+
+#pragma omp requires unified_shared_memory
+
+int a[10];
+extern void foo (void);
+
+int
+main (void)
+{
+ #pragma omp target
+ for (int i = 0; i < 10; i++)
+ a[i] = 0;
+
+ foo ();
+ return 0;
+}
+
+/* { dg-error "OpenMP 'requires' directive with non-identical clauses in multiple compilation units: 'unified_shared_memory' vs. 'unified_address'" "" { target *-*-* } 0 } */
+/* { dg-excess-errors "Ignore messages like: errors during merging of translation units|mkoffload returned 1 exit status" } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-2-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-2-aux.c
new file mode 100644
index 0000000000000000000000000000000000000000..617577448ed12024568336d07a3feddede9da5fb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-2-aux.c
@@ -0,0 +1,9 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+int x;
+
+void foo (void)
+{
+ #pragma omp target
+ x = 1;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-2.c b/libgomp/testsuite/libgomp.c-c++-common/requires-2.c
new file mode 100644
index 0000000000000000000000000000000000000000..be1830d0c46eef9c75a3385ba17705785c751faf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-2.c
@@ -0,0 +1,25 @@
+/* { dg-do link { target offloading_enabled } } */
+/* { dg-additional-options "-foffload=disable -flto" } */
+/* { dg-additional-sources requires-2-aux.c } */
+
+/* Check diagnostic by host's lto1.
+ Other file does not have any 'omp requires'. */
+
+#pragma omp requires unified_shared_memory
+
+int a[10];
+extern void foo (void);
+
+int
+main (void)
+{
+ #pragma omp target
+ for (int i = 0; i < 10; i++)
+ a[i] = 0;
+
+ foo ();
+ return 0;
+}
+
+/* { dg-error "OpenMP 'requires' directive with 'unified_shared_memory' specified only in some compilation units" "" { target *-*-* } 0 } */
+/* { dg-excess-errors "Ignore messages like: errors during merging of translation units|mkoffload returned 1 exit status" } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-3-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-3-aux.c
new file mode 100644
index 0000000000000000000000000000000000000000..bdca662e42f9043224d02d1639e18b3428116b8e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-3-aux.c
@@ -0,0 +1,11 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires unified_address
+
+int x;
+
+void foo (void)
+{
+ #pragma omp target
+ x = 1;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-3.c b/libgomp/testsuite/libgomp.c-c++-common/requires-3.c
new file mode 100644
index 0000000000000000000000000000000000000000..4b07ffdd09b42fc3f0c24a5750b3a0593cd59132
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-3.c
@@ -0,0 +1,24 @@
+/* { dg-do link { target offloading_enabled } } */
+/* { dg-additional-sources requires-3-aux.c } */
+
+/* Check diagnostic by device-compiler's lto1.
+ Other file uses: 'requires unified_address'. */
+
+#pragma omp requires unified_address,unified_shared_memory
+
+int a[10];
+extern void foo (void);
+
+int
+main (void)
+{
+ #pragma omp target
+ for (int i = 0; i < 10; i++)
+ a[i] = 0;
+
+ foo ();
+ return 0;
+}
+
+/* { dg-error "OpenMP 'requires' directive with non-identical clauses in multiple compilation units: 'unified_address, unified_shared_memory' vs. 'unified_address'" "" { target *-*-* } 0 } */
+/* { dg-excess-errors "Ignore messages like: errors during merging of translation units|mkoffload returned 1 exit status" } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-4-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-4-aux.c
new file mode 100644
index 0000000000000000000000000000000000000000..b8b51ae8ca71ef9c69687064a5a3b5d7f9592998
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-4-aux.c
@@ -0,0 +1,13 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires reverse_offload
+
+/* Note: The file does not have neither of:
+ declare target directives, device constructs or device routines. */
+
+int x;
+
+void foo (void)
+{
+ x = 1;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-4.c b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
new file mode 100644
index 0000000000000000000000000000000000000000..128fdbb8463d640faf76f1a31cc04fc57ceb54ec
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
@@ -0,0 +1,23 @@
+/* { dg-do link { target offloading_enabled } } */
+/* { dg-additional-options "-flto" } */
+/* { dg-additional-sources requires-4-aux.c } */
+
+/* Check diagnostic by device-compiler's or host compiler's lto1.
+ Other file uses: 'requires reverse_offload', but that's inactive as
+ there are no declare target directives, device constructs nor device routines */
+
+#pragma omp requires unified_address,unified_shared_memory
+
+int a[10];
+extern void foo (void);
+
+int
+main (void)
+{
+ #pragma omp target
+ for (int i = 0; i < 10; i++)
+ a[i] = 0;
+
+ foo ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-5-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-5-aux.c
new file mode 100644
index 0000000000000000000000000000000000000000..d223749f0a12f4eeca374d46feed07ea05f61751
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-5-aux.c
@@ -0,0 +1,11 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires unified_shared_memory, unified_address, reverse_offload
+
+int x;
+
+void foo (void)
+{
+ #pragma omp target
+ x = 1;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-5.c b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
new file mode 100644
index 0000000000000000000000000000000000000000..c1e5540cfc51c2b6b46dafb5b41a022a0fc41df5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
@@ -0,0 +1,21 @@
+/* { dg-do run { target { offload_target_nvptx || offload_target_amdgcn } } } */
+/* { dg-additional-sources requires-5-aux.c } */
+
+#pragma omp requires unified_shared_memory, unified_address, reverse_offload
+
+int a[10];
+extern void foo (void);
+
+int
+main (void)
+{
+ #pragma omp target
+ for (int i = 0; i < 10; i++)
+ a[i] = 0;
+
+ foo ();
+ return 0;
+}
+
+/* (Only) if GOMP_DEBUG=1, should print at runtime the following:
+ "devices present but 'omp requires unified_address, unified_shared_memory, reverse_offload' cannot be fulfilled" */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-6.c b/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
new file mode 100644
index 0000000000000000000000000000000000000000..b00c7459bbc3094ccfa704cf0fc4a19695f4f4b5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
@@ -0,0 +1,17 @@
+#pragma omp requires unified_shared_memory, unified_address, reverse_offload
+
+/* The requires line is not active as there is none of:
+ declare target directives, device constructs or device routines.
+ Thus, this code is expected to work everywhere. */
+
+int a[10];
+extern void foo (void);
+
+int
+main (void)
+{
+ for (int i = 0; i < 10; i++)
+ a[i] = 0;
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-7-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-7-aux.c
new file mode 100644
index 0000000000000000000000000000000000000000..0916db8a0ce07a3dd779748fa3fd865307d25e49
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-7-aux.c
@@ -0,0 +1,11 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires unified_address
+
+int x;
+
+void foo (void)
+{
+ x = 1;
+ #pragma omp target enter data map(always,to: x)
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-7.c b/libgomp/testsuite/libgomp.c-c++-common/requires-7.c
new file mode 100644
index 0000000000000000000000000000000000000000..c94a4c1084636029a2d39f387dfb92976e1d60e7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-7.c
@@ -0,0 +1,24 @@
+/* { dg-do link { target { offload_target_nvptx || offload_target_amdgcn } } } */
+/* { dg-additional-sources requires-7-aux.c } */
+
+/* Check diagnostic by device-compiler's lto1.
+ Other file uses: 'requires unified_address'. */
+
+#pragma omp requires unified_shared_memory
+
+int a[10];
+extern void foo (void);
+
+int
+main (void)
+{
+ #pragma omp target
+ for (int i = 0; i < 10; i++)
+ a[i] = 0;
+
+ foo ();
+ return 0;
+}
+
+/* { dg-error "OpenMP 'requires' directive with non-identical clauses in multiple compilation units: 'unified_shared_memory' vs. 'unified_address'" "" { target *-*-* } 0 } */
+/* { dg-excess-errors "Ignore messages like: errors during merging of translation units|mkoffload returned 1 exit status" } */
diff --git a/libgomp/testsuite/libgomp.fortran/requires-1-aux.f90 b/libgomp/testsuite/libgomp.fortran/requires-1-aux.f90
new file mode 100644
index 0000000000000000000000000000000000000000..a18caeb4c694545a31783b3707cc4a59e6173719
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/requires-1-aux.f90
@@ -0,0 +1,14 @@
+! { dg-skip-if "" { *-*-* } }
+
+module m
+ integer x
+end module m
+
+subroutine foo
+ use m
+ implicit none
+ !$omp requires unified_address
+
+ x = 1
+ !$omp target enter data map(always,to: x)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/requires-1.f90 b/libgomp/testsuite/libgomp.fortran/requires-1.f90
new file mode 100644
index 0000000000000000000000000000000000000000..33741af15f1a0ab8ded42900f9fe1fd026c6f52a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/requires-1.f90
@@ -0,0 +1,26 @@
+! { dg-do link { target { offload_target_nvptx || offload_target_amdgcn } } }
+! { dg-additional-sources requires-1-aux.f90 }
+
+! Check diagnostic by device-compiler's lto1.
+! Other file uses: 'requires unified_address'.
+
+module m
+ integer :: a(10)
+ interface
+ subroutine foo
+ end
+ end interface
+end
+
+program main
+ !$omp requires unified_shared_memory
+
+ !$omp target
+ a = 0
+ !$omp end target
+
+ call foo ()
+end
+
+! { dg-error "OpenMP 'requires' directive with non-identical clauses in multiple compilation units: 'unified_shared_memory' vs. 'unified_address'" "" { target *-*-* } 0 }
+! { dg-excess-errors "Ignore messages like: errors during merging of translation units|mkoffload returned 1 exit status" }
diff --git a/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
index d1678d0514e91e775fff7fe361b861f706b53782..33bae0650b41c6fffe0036f8b1bf4c0ddc531c64 100644
--- a/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
+++ b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
@@ -168,8 +168,12 @@ GOMP_OFFLOAD_get_type (void)
}
extern "C" int
-GOMP_OFFLOAD_get_num_devices (void)
+GOMP_OFFLOAD_get_num_devices (unsigned int omp_requires_mask)
{
+ /* Return -1 if no omp_requires_mask cannot be fulfilled but
+ devices were present. */
+ if (num_devices > 0 && omp_requires_mask != 0)
+ return -1;
TRACE ("(): return %d", num_devices);
return num_devices;
}