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");
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;
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");
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");
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)
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);
/* 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);
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"
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)
if (!in)
fatal_error (input_location, "cannot open intermediate gcn obj file");
- process_obj (in, cfile);
+ process_obj (in, cfile, omp_requires);
fclose (in);
}
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);
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;)
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"
" 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"
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);
}
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);
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);
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);
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;
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)
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;
}
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:
/* 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;
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;
#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"
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
(*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);
}
}
+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. */
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++]))
{
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)
LTO mode. */
if (do_force_output)
cgraph_node::get (fn_decl)->mark_force_output ();
+ tmp_decl = fn_decl;
}
else if (tag == LTO_symtab_variable)
{
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,
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. */
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;
#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" } */
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)
/* { 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)
;
#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)
;
/* { 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)
/* 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" } */
;
}
-#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 ()
+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
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
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
!$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
--- /dev/null
+! { 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
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)
n = omp_get_thread_num ()
!$omp end target
-end
\ No newline at end of file
+end
! 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" } }
!
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)
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
end module m
subroutine foo()
- !$omp requires reverse_offload ! { dg-error "REQUIRES directive is not yet supported" }
+ !$omp requires reverse_offload
block
block
block
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)
/* 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
#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. */
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);
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
@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
@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
}
static int
-host_get_num_devices (void)
+host_get_num_devices (unsigned int omp_requires_mask __attribute__((unused)))
{
return 1;
}
{
union goacc_property_value nullval = { .val = 0 };
- if (n >= host_get_num_devices ())
+ if (n >= host_get_num_devices (0))
return nullval;
switch (prop)
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)
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)
{
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);
/* 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++)
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);
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;
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);
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);
/* 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;
}
}
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
# include <inttypes.h> /* For PRIu64. */
#endif
#include <string.h>
+#include <stdio.h> /* For snprintf. */
#include <assert.h>
#include <errno.h>
/* 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 *
}
}
+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. */
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++)
{
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. */
--- /dev/null
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires unified_address
+
+int x;
+
+void foo (void)
+{
+ #pragma omp target
+ x = 1;
+}
--- /dev/null
+/* { 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" } */
--- /dev/null
+/* { dg-skip-if "" { *-*-* } } */
+
+int x;
+
+void foo (void)
+{
+ #pragma omp target
+ x = 1;
+}
--- /dev/null
+/* { 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" } */
--- /dev/null
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires unified_address
+
+int x;
+
+void foo (void)
+{
+ #pragma omp target
+ x = 1;
+}
--- /dev/null
+/* { 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" } */
--- /dev/null
+/* { 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;
+}
--- /dev/null
+/* { 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;
+}
--- /dev/null
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires unified_shared_memory, unified_address, reverse_offload
+
+int x;
+
+void foo (void)
+{
+ #pragma omp target
+ x = 1;
+}
--- /dev/null
+/* { 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" */
--- /dev/null
+#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;
+}
--- /dev/null
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires unified_address
+
+int x;
+
+void foo (void)
+{
+ x = 1;
+ #pragma omp target enter data map(always,to: x)
+}
--- /dev/null
+/* { 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" } */
--- /dev/null
+! { 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
--- /dev/null
+! { 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" }
}
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;
}
projects / platform / upstream / gcc.git / commitdiff