OpenMP: Move omp requires checks to libgomp Handle reverse_offload, unified_address, and unified_shared_memory requirements in libgomp by putting them into the .gnu.gomp_requires section. For all in-principle supported devices, if a requirement cannot be fulfilled, the device is excluded from the (supported) devices list. Currently, none of those requirements are marked as supported for any of the non-host devices. Additionally, libgomp checks for consistency across the entire .gnu.gomp_requires section, matching the requirements set by the OpenMP spec. gcc/c/ChangeLog: * c-parser.cc (c_parser_declaration_or_fndef): Set OMP_REQUIRES_TARGET_USED in omp_requires_mask if function has "omp declare target" attribute. (c_parser_omp_target_data): Set OMP_REQUIRES_TARGET_USED in omp_requires_mask. (c_parser_omp_target_enter_data): Likewise. (c_parser_omp_target_exit_data): Likewise. (c_parser_omp_requires): Remove sorry. gcc/cp/ChangeLog: * parser.cc (cp_parser_simple_declaration): Set OMP_REQUIRES_TARGET_USED in omp_requires_mask if function has "omp declare target" attribute. (cp_parser_omp_target_data): Set OMP_REQUIRES_TARGET_USED in omp_requires_mask. (cp_parser_omp_target_enter_data): Likewise. (cp_parser_omp_target_exit_data): Likewise. (cp_parser_omp_requires): Remove sorry. gcc/fortran/ChangeLog: * openmp.cc (gfc_match_omp_requires): Remove "not implemented yet". * parse.cc: Include "tree.h" and "omp-general.h". (gfc_parse_file): Add code to merge omp_requires to omp_requires_mask. gcc/ChangeLog: * omp-general.h (omp_runtime_api_call): New prototype. * omp-general.cc (omp_runtime_api_call): Added device_api_only arg and moved from ... * omp-low.cc (omp_runtime_api_call): ... here. (scan_omp_1_stmt): Update call. * gimplify.cc (gimplify_call_expr): Call omp_runtime_api_call. * omp-offload.cc (omp_finish_file): Add code to create OpenMP requires mask variable in .gnu.gomp_requires section, if needed. include/ChangeLog: * gomp-constants.h (GOMP_REQUIRES_UNIFIED_ADDRESS, GOMP_REQUIRES_UNIFIED_SHARED_MEMORY, GOMP_REQUIRES_REVERSE_OFFLOAD): New. libgcc/ChangeLog: * offloadstuff.c (__requires_mask_table, __requires_mask_table_end): New symbols to mark start and end of the .gnu.gomp_requires section. libgomp/ChangeLog: * libgomp-plugin.h (GOMP_OFFLOAD_get_num_devices): Add omp_requires_mask arg. * plugin/plugin-gcn.c (GOMP_OFFLOAD_get_num_devices): Likewise; return -1 when device available but omp_requires_mask != 0. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_get_num_devices): Likewise. * oacc-host.c (host_get_num_devices, host_openacc_get_property): Update call. * oacc-init.c (resolve_device, acc_init_1, acc_shutdown_1, goacc_attach_host_thread_to_device, acc_get_num_devices, acc_set_device_num, get_property_any): Likewise. * target.c: (__requires_mask_table, __requires_mask_table_end): Declare weak extern symbols. (gomp_requires_to_name): New. (gomp_target_init): Add code to check .gnu._gomp_requires section mask values for inconsistencies; warn when requirements makes an existing device unsupported. * testsuite/libgomp.c-c++-common/requires-1-aux.c: New test. * testsuite/libgomp.c-c++-common/requires-1.c: New test. * testsuite/libgomp.c-c++-common/requires-2-aux.c: New test. * testsuite/libgomp.c-c++-common/requires-2.c: New test. liboffloadmic/ChangeLog: * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_num_devices): Return -1 when device available but omp_requires_mask != 0. gcc/testsuite/ChangeLog: * c-c++-common/gomp/requires-4.c: Update dg-*. * c-c++-common/gomp/target-device-ancestor-2.c: Likewise. * c-c++-common/gomp/target-device-ancestor-3.c: Likewise. * c-c++-common/gomp/target-device-ancestor-4.c: Likewise. * c-c++-common/gomp/target-device-ancestor-5.c: Likewise. * gfortran.dg/gomp/target-device-ancestor-3.f90: Likewise. * gfortran.dg/gomp/target-device-ancestor-4.f90: Likewise. * gfortran.dg/gomp/target-device-ancestor-2.f90: Likewise. Move post-FE checks to ... * gfortran.dg/gomp/target-device-ancestor-2a.f90: ... this new file. Co-authored-by: Chung-Lin Tang Co-authored-by: Thomas Schwinge gcc/c/c-parser.cc | 21 +++- gcc/cp/parser.cc | 20 ++- gcc/fortran/openmp.cc | 4 - gcc/fortran/parse.cc | 21 ++++ gcc/gimplify.cc | 3 + gcc/omp-general.cc | 137 +++++++++++++++++++++ gcc/omp-general.h | 1 + gcc/omp-low.cc | 135 +------------------- gcc/omp-offload.cc | 29 +++++ gcc/testsuite/c-c++-common/gomp/requires-4.c | 2 - .../c-c++-common/gomp/target-device-ancestor-2.c | 10 +- .../c-c++-common/gomp/target-device-ancestor-3.c | 2 +- .../c-c++-common/gomp/target-device-ancestor-4.c | 4 +- .../c-c++-common/gomp/target-device-ancestor-5.c | 2 +- .../gfortran.dg/gomp/target-device-ancestor-2.f90 | 70 +---------- .../gfortran.dg/gomp/target-device-ancestor-2a.f90 | 80 ++++++++++++ .../gfortran.dg/gomp/target-device-ancestor-3.f90 | 6 +- .../gfortran.dg/gomp/target-device-ancestor-4.f90 | 6 +- include/gomp-constants.h | 6 + libgcc/offloadstuff.c | 9 ++ libgomp/libgomp-plugin.h | 2 +- libgomp/oacc-host.c | 4 +- libgomp/oacc-init.c | 16 +-- libgomp/plugin/plugin-gcn.c | 6 +- libgomp/plugin/plugin-nvptx.c | 9 +- libgomp/target.c | 66 +++++++++- .../libgomp.c-c++-common/requires-1-aux.c | 11 ++ .../testsuite/libgomp.c-c++-common/requires-1.c | 21 ++++ .../libgomp.c-c++-common/requires-2-aux.c | 11 ++ .../testsuite/libgomp.c-c++-common/requires-2.c | 20 +++ liboffloadmic/plugin/libgomp-plugin-intelmic.cpp | 6 +- 31 files changed, 499 insertions(+), 241 deletions(-) diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 1704a52be12..4748ce04737 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -2488,6 +2488,12 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok, break; } + if (flag_openmp + && lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (current_function_decl))) + omp_requires_mask + = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED); + if (DECL_DECLARED_INLINE_P (current_function_decl)) tv = TV_PARSE_INLINE; else @@ -20915,6 +20921,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"); @@ -21057,6 +21067,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"); @@ -21143,6 +21157,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"); @@ -22763,9 +22781,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 % directive not " - "supported yet", p); if (p) c_parser_consume_token (parser); if (this_req) diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index da2f370cdca..6e26d123370 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -15389,6 +15389,11 @@ cp_parser_simple_declaration (cp_parser* parser, /* Otherwise, we're done with the list of declarators. */ else { + if (flag_openmp && lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (decl))) + omp_requires_mask + = (enum omp_requires) (omp_requires_mask + | OMP_REQUIRES_TARGET_USED); pop_deferring_access_checks (); cp_finalize_omp_declare_simd (parser, &odsd); return; @@ -44287,6 +44292,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); @@ -44390,6 +44399,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); @@ -44481,6 +44494,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); @@ -46861,9 +46878,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 % 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 d12cec43d64..7790ef34664 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -5481,10 +5481,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 7356d1b5a3a..b142e169a5c 100644 --- a/gcc/fortran/parse.cc +++ b/gcc/fortran/parse.cc @@ -6908,6 +6908,27 @@ done: break; } + if (omp_requires & OMP_REQ_TARGET_MASK) + { + 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/gimplify.cc b/gcc/gimplify.cc index cd1796643d7..3fe4571d677 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -3644,6 +3644,9 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value) return GS_OK; } } + if (fndecl && flag_openmp && omp_runtime_api_call (fndecl, true)) + omp_requires_mask + = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED); /* Remember the original function pointer type. */ fnptrtype = TREE_TYPE (CALL_EXPR_FN (*expr_p)); diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc index a406c578f33..120bcaa10b2 100644 --- a/gcc/omp-general.cc +++ b/gcc/omp-general.cc @@ -89,6 +89,143 @@ omp_privatize_by_reference (tree decl) return lang_hooks.decls.omp_privatize_by_reference (decl); } +/* Return true if FNDECL is an omp_* runtime API call; with device_api_only set, + returns true only for device API calls. */ + +bool +omp_runtime_api_call (const_tree fndecl, bool device_api_only) +{ + tree declname = DECL_NAME (fndecl); + if (!declname + || (DECL_CONTEXT (fndecl) != NULL_TREE + && TREE_CODE (DECL_CONTEXT (fndecl)) != TRANSLATION_UNIT_DECL) + || !TREE_PUBLIC (fndecl)) + return false; + + const char *name = IDENTIFIER_POINTER (declname); + if (!startswith (name, "omp_")) + return false; + + static const char *omp_runtime_apis[] = + { + /* This array has 6 sections. First omp_* calls that don't + have any suffixes and are non-device APIs. */ + "aligned_alloc", + "aligned_calloc", + "alloc", + "calloc", + "free", + "realloc", + NULL, + /* Now likewise but for device API. */ + "get_mapped_ptr", + "target_alloc", + "target_associate_ptr", + "target_disassociate_ptr", + "target_free", + "target_is_accessible", + "target_is_present", + "target_memcpy", + "target_memcpy_async", + "target_memcpy_rect", + "target_memcpy_rect_async", + NULL, + /* Now omp_* calls that are available as omp_* and omp_*_; however, the + DECL_NAME is always omp_* without tailing underscore. Non device. */ + "capture_affinity", + "destroy_allocator", + "destroy_lock", + "destroy_nest_lock", + "display_affinity", + "fulfill_event", + "get_active_level", + "get_affinity_format", + "get_cancellation", + "get_default_allocator", + "get_default_device", + "get_dynamic", + "get_level", + "get_max_active_levels", + "get_max_task_priority", + "get_max_teams", + "get_max_threads", + "get_nested", + "get_num_devices", + "get_num_places", + "get_num_procs", + "get_num_teams", + "get_num_threads", + "get_partition_num_places", + "get_place_num", + "get_proc_bind", + "get_supported_active_levels", + "get_team_num", + "get_teams_thread_limit", + "get_thread_limit", + "get_thread_num", + "get_wtick", + "get_wtime", + "in_final", + "in_parallel", + "init_lock", + "init_nest_lock", + "pause_resource", + "pause_resource_all", + "set_affinity_format", + "set_default_allocator", + "set_lock", + "set_nest_lock", + "test_lock", + "test_nest_lock", + "unset_lock", + "unset_nest_lock", + NULL, + /* And device APIs. */ + "get_device_num", + "get_initial_device", + "is_initial_device", /* Even if it does not require init'ed devices. */ + NULL, + /* And finally calls available as omp_*, omp_*_ and omp_*_8_; however, + as DECL_NAME only omp_* and omp_*_8 appear. For non device. */ + "display_env", + "get_ancestor_thread_num", + "init_allocator", + "get_partition_place_nums", + "get_place_num_procs", + "get_place_proc_ids", + "get_schedule", + "get_team_size", + "set_default_device", + "set_dynamic", + "set_max_active_levels", + "set_nested", + "set_num_teams", + "set_num_threads", + "set_schedule", + "set_teams_thread_limit", + NULL, + /* And for device APIs. (Currently none.) */ + }; + + int mode = 0; + for (unsigned i = 0; i < ARRAY_SIZE (omp_runtime_apis); i++) + { + if (omp_runtime_apis[i] == NULL) + { + mode++; + continue; + } + if (device_api_only && mode % 2 != 0) + continue; + size_t len = strlen (omp_runtime_apis[i]); + if (strncmp (name + 4, omp_runtime_apis[i], len) == 0 + && (name[4 + len] == '\0' + || (mode > 1 && strcmp (name + 4 + len, "_8") == 0))) + return true; + } + return false; +} + /* Adjust *COND_CODE and *N2 so that the former is either LT_EXPR or GT_EXPR, given that V is the loop index variable and STEP is loop step. */ diff --git a/gcc/omp-general.h b/gcc/omp-general.h index 7a94831e8f5..f1be9f23ef7 100644 --- a/gcc/omp-general.h +++ b/gcc/omp-general.h @@ -95,6 +95,7 @@ extern tree omp_find_clause (tree clauses, enum omp_clause_code kind); extern bool omp_is_allocatable_or_ptr (tree decl); extern tree omp_check_optional_argument (tree decl, bool for_present_check); extern bool omp_privatize_by_reference (tree decl); +extern bool omp_runtime_api_call (const_tree fndecl, bool device_api_only); extern void omp_adjust_for_condition (location_t loc, enum tree_code *cond_code, tree *n2, tree v, tree step); extern tree omp_get_for_step_from_incr (location_t loc, tree incr); diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index f976e3a1549..243fa27a62f 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -3989,134 +3989,6 @@ setjmp_or_longjmp_p (const_tree fndecl) return !strcmp (name, "setjmp") || !strcmp (name, "longjmp"); } -/* Return true if FNDECL is an omp_* runtime API call. */ - -static bool -omp_runtime_api_call (const_tree fndecl) -{ - tree declname = DECL_NAME (fndecl); - if (!declname - || (DECL_CONTEXT (fndecl) != NULL_TREE - && TREE_CODE (DECL_CONTEXT (fndecl)) != TRANSLATION_UNIT_DECL) - || !TREE_PUBLIC (fndecl)) - return false; - - const char *name = IDENTIFIER_POINTER (declname); - if (!startswith (name, "omp_")) - return false; - - static const char *omp_runtime_apis[] = - { - /* This array has 3 sections. First omp_* calls that don't - have any suffixes. */ - "aligned_alloc", - "aligned_calloc", - "alloc", - "calloc", - "free", - "get_mapped_ptr", - "realloc", - "target_alloc", - "target_associate_ptr", - "target_disassociate_ptr", - "target_free", - "target_is_accessible", - "target_is_present", - "target_memcpy", - "target_memcpy_async", - "target_memcpy_rect", - "target_memcpy_rect_async", - NULL, - /* Now omp_* calls that are available as omp_* and omp_*_; however, the - DECL_NAME is always omp_* without tailing underscore. */ - "capture_affinity", - "destroy_allocator", - "destroy_lock", - "destroy_nest_lock", - "display_affinity", - "fulfill_event", - "get_active_level", - "get_affinity_format", - "get_cancellation", - "get_default_allocator", - "get_default_device", - "get_device_num", - "get_dynamic", - "get_initial_device", - "get_level", - "get_max_active_levels", - "get_max_task_priority", - "get_max_teams", - "get_max_threads", - "get_nested", - "get_num_devices", - "get_num_places", - "get_num_procs", - "get_num_teams", - "get_num_threads", - "get_partition_num_places", - "get_place_num", - "get_proc_bind", - "get_supported_active_levels", - "get_team_num", - "get_teams_thread_limit", - "get_thread_limit", - "get_thread_num", - "get_wtick", - "get_wtime", - "in_final", - "in_parallel", - "init_lock", - "init_nest_lock", - "is_initial_device", - "pause_resource", - "pause_resource_all", - "set_affinity_format", - "set_default_allocator", - "set_lock", - "set_nest_lock", - "test_lock", - "test_nest_lock", - "unset_lock", - "unset_nest_lock", - NULL, - /* And finally calls available as omp_*, omp_*_ and omp_*_8_; however, - as DECL_NAME only omp_* and omp_*_8 appear. */ - "display_env", - "get_ancestor_thread_num", - "init_allocator", - "get_partition_place_nums", - "get_place_num_procs", - "get_place_proc_ids", - "get_schedule", - "get_team_size", - "set_default_device", - "set_dynamic", - "set_max_active_levels", - "set_nested", - "set_num_teams", - "set_num_threads", - "set_schedule", - "set_teams_thread_limit" - }; - - int mode = 0; - for (unsigned i = 0; i < ARRAY_SIZE (omp_runtime_apis); i++) - { - if (omp_runtime_apis[i] == NULL) - { - mode++; - continue; - } - size_t len = strlen (omp_runtime_apis[i]); - if (strncmp (name + 4, omp_runtime_apis[i], len) == 0 - && (name[4 + len] == '\0' - || (mode > 1 && strcmp (name + 4 + len, "_8") == 0))) - return true; - } - return false; -} - /* Helper function for scan_omp. Callback for walk_gimple_stmt used to scan for OMP directives in @@ -4171,7 +4043,8 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, omp_context *octx = ctx; if (gimple_code (ctx->stmt) == GIMPLE_OMP_SCAN && ctx->outer) octx = ctx->outer; - if (octx->order_concurrent && omp_runtime_api_call (fndecl)) + if (octx->order_concurrent + && omp_runtime_api_call (fndecl, false)) { remove = true; error_at (gimple_location (stmt), @@ -4179,7 +4052,7 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, "% clause", fndecl); } if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS - && omp_runtime_api_call (fndecl) + && omp_runtime_api_call (fndecl, false) && ((IDENTIFIER_LENGTH (DECL_NAME (fndecl)) != strlen ("omp_get_num_teams")) || strcmp (IDENTIFIER_POINTER (DECL_NAME (fndecl)), @@ -4197,7 +4070,7 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET && (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_REGION) - && omp_runtime_api_call (fndecl)) + && omp_runtime_api_call (fndecl, false)) { tree tgt_clauses = gimple_omp_target_clauses (ctx->stmt); tree c = omp_find_clause (tgt_clauses, OMP_CLAUSE_DEVICE); diff --git a/gcc/omp-offload.cc b/gcc/omp-offload.cc index ad4e772015e..998abab0f11 100644 --- a/gcc/omp-offload.cc +++ b/gcc/omp-offload.cc @@ -397,6 +397,27 @@ omp_finish_file (void) unsigned num_funcs = vec_safe_length (offload_funcs); unsigned num_vars = vec_safe_length (offload_vars); + if (flag_openmp && (omp_requires_mask & OMP_REQUIRES_TARGET_USED) != 0) + { + if (targetm_common.have_named_sections) + { + const char *requires_section = ".gnu.gomp_requires"; + tree maskvar = build_decl (UNKNOWN_LOCATION, VAR_DECL, + get_identifier (".gomp_requires_mask"), + unsigned_type_node); + SET_DECL_ALIGN (maskvar, TYPE_ALIGN (unsigned_type_node)); + TREE_STATIC (maskvar) = 1; + DECL_INITIAL (maskvar) + = build_int_cst (unsigned_type_node, + ((unsigned int) omp_requires_mask + & (OMP_REQUIRES_UNIFIED_ADDRESS + | OMP_REQUIRES_UNIFIED_SHARED_MEMORY + | OMP_REQUIRES_REVERSE_OFFLOAD))); + set_decl_section_name (maskvar, requires_section); + varpool_node::finalize_decl (maskvar); + } + } + if (num_funcs == 0 && num_vars == 0) return; @@ -442,6 +463,14 @@ omp_finish_file (void) } else { +#ifndef ACCEL_COMPILER + if (flag_openmp + && (omp_requires_mask & OMP_REQUIRES_TARGET_USED) + && (omp_requires_mask & (OMP_REQUIRES_UNIFIED_ADDRESS + | OMP_REQUIRES_UNIFIED_SHARED_MEMORY + | OMP_REQUIRES_REVERSE_OFFLOAD))) + sorry ("OpenMP device offloading is not supported for this target"); +#endif for (unsigned i = 0; i < num_funcs; i++) { tree it = (*offload_funcs)[i]; diff --git a/gcc/testsuite/c-c++-common/gomp/requires-4.c b/gcc/testsuite/c-c++-common/gomp/requires-4.c index 88ba7746cf8..8f45d83ea6e 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/target-device-ancestor-2.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c index cf05c505004..b16e701bd5a 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 ea6e5a0cf6c..d16590107d2 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 b4b5620bbc0..241234f8daf 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 b6ff84bcdab..b1520ff0636 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/target-device-ancestor-2.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2.f90 index 117a1d000a5..230c690d84c 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 00000000000..feb76fe2144 --- /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 f1145bde2ec..e8975e6a08b 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 63872fa51fb..ab56e2d1d52 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/include/gomp-constants.h b/include/gomp-constants.h index 701d33dae49..ebf6978b697 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -330,6 +330,12 @@ 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 + /* HSA specific data structures. */ /* Identifiers of device-specific target arguments. */ diff --git a/libgcc/offloadstuff.c b/libgcc/offloadstuff.c index 10e1fe19c8e..b2282924fb4 100644 --- a/libgcc/offloadstuff.c +++ b/libgcc/offloadstuff.c @@ -54,6 +54,9 @@ const void *const __offload_var_table[0] __attribute__ ((__used__, visibility ("hidden"), section (OFFLOAD_VAR_TABLE_SECTION_NAME))) = { }; +const unsigned int const __requires_mask_table[0] + __attribute__ ((__used__, section (".gnu.gomp_requires"))) = { }; + #elif defined CRT_END const void *const __offload_funcs_end[0] @@ -63,6 +66,9 @@ const void *const __offload_vars_end[0] __attribute__ ((__used__, visibility ("hidden"), section (OFFLOAD_VAR_TABLE_SECTION_NAME))) = { }; +const unsigned int const __requires_mask_table_end[0] + __attribute__ ((__used__, section (".gnu.gomp_requires"))) = { }; + #elif defined CRT_TABLE extern const void *const __offload_func_table[]; @@ -77,6 +83,9 @@ const void *const __OFFLOAD_TABLE__[] &__offload_var_table, &__offload_vars_end }; +extern const unsigned int const __requires_mask_table[]; +extern const unsigned int const __requires_mask_table_end[]; + #else /* ! CRT_BEGIN && ! CRT_END && ! CRT_TABLE */ #error "One of CRT_BEGIN, CRT_END or CRT_TABLE must be defined." #endif diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h index 07ab700b80c..ab3ed638475 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/oacc-host.c b/libgomp/oacc-host.c index 5bb889926d3..eb11b9cf16a 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 1565aa0f290..42c3e74e6ba 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 1c0436842da..ea327bf2ca0 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 387bcbbc52a..bc63e274cdf 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 4740f8a45d3..0fd3f7f47ad 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -36,6 +36,7 @@ # include /* For PRIu64. */ #endif #include +#include /* For snprintf. */ #include #include @@ -98,6 +99,13 @@ static int num_devices; /* Number of GOMP_OFFLOAD_CAP_OPENMP_400 devices. */ static int num_devices_openmp; +/* Start/end of .gnu.gomp.requires section of program, defined in + crtoffloadbegin/end.o. */ +__attribute__((weak)) +extern const unsigned int __requires_mask_table[]; +__attribute__((weak)) +extern const unsigned int __requires_mask_table_end[]; + /* Similar to gomp_realloc, but release register_lock before gomp_fatal. */ static void * @@ -4085,6 +4093,20 @@ gomp_target_fini (void) } } +static void +gomp_requires_to_name (char *buf, size_t size, unsigned 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 initializes the runtime for offloading. It parses the list of offload plugins, and tries to load these. On return, the variables NUM_DEVICES and NUM_DEVICES_OPENMP @@ -4106,6 +4128,35 @@ gomp_target_init (void) if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_DISABLED) return; + /* Mask of requires directive clause values, summarized from + .gnu.gomp.requires section. Offload plugins are queried with this mask to see + if all required features are supported. */ + unsigned int requires_mask = 0; + const unsigned int *mask_ptr = __requires_mask_table; + bool error_emitted = false; + while (mask_ptr != __requires_mask_table_end) + { + if (requires_mask == 0) + requires_mask = *mask_ptr; + else if (requires_mask != *mask_ptr) + { + if (!error_emitted) + { + char buf[64], buf2[64]; + gomp_requires_to_name (buf, sizeof (buf), requires_mask); + gomp_requires_to_name (buf2, sizeof (buf2), *mask_ptr); + gomp_error ("requires-directive clause inconsistency between " + "compilation units detected: '%s' vs. '%s'", + buf, buf2); + error_emitted = true; + } + /* This is inconsistent, but still merge to query for all features + later. */ + requires_mask |= *mask_ptr; + } + mask_ptr++; + } + cur = OFFLOAD_PLUGINS; if (*cur) do @@ -4132,8 +4183,19 @@ 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) + new_num_devs = current_device.get_num_devices_func (requires_mask); + if (new_num_devs < 0) + { + char buf[64]; + gomp_requires_to_name (buf, sizeof (buf), requires_mask); + char *name = (char *) malloc (cur_len + 1); + memcpy (name, cur, cur_len); + name[cur_len] = '\0'; + GOMP_PLUGIN_error ("note: %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 00000000000..8b9341523c6 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-1-aux.c @@ -0,0 +1,11 @@ +/* { dg-skip-if "" { *-*-* } } */ + +#pragma omp requires reverse_offload + +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 00000000000..990b4e9817d --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-1.c @@ -0,0 +1,21 @@ +/* { dg-skip-if "" { ! offloading_enabled } } */ +/* { dg-additional-sources requires-1-aux.c } */ + +#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-output "libgomp: requires-directive clause inconsistency between compilation units detected" } */ +/* { dg-prune-output "nvptx device present but 'omp requires unified_shared_memory, reverse_offload, reverse_offload' cannot be fulfilled" } */ 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 00000000000..4077648347d --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-2-aux.c @@ -0,0 +1,11 @@ +/* { dg-skip-if "" { *-*-* } } */ + +#pragma omp requires unified_shared_memory + +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 00000000000..bc55ab001e9 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-2.c @@ -0,0 +1,20 @@ +/* { dg-additional-sources requires-2-aux.c } */ +/* { dg-require-effective-target offload_device } */ + +#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-output "devices present but 'omp requires unified_shared_memory' cannot be fulfilled" } */ diff --git a/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp index d1678d0514e..33bae0650b4 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; }