diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index c77d9fccdc2..e685b26746e 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -2475,6 +2475,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 @@ -19556,6 +19562,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"); @@ -19698,6 +19708,10 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser, return NULL_TREE; } + 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"); @@ -19784,6 +19798,10 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser, return NULL_TREE; } + 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"); @@ -21371,7 +21389,7 @@ c_parser_omp_requires (c_parser *parser) c_parser_skip_to_pragma_eol (parser, false); return; } - if (p) + if (this_req == OMP_REQUIRES_DYNAMIC_ALLOCATORS) sorry_at (cloc, "%qs clause on % directive not " "supported yet", p); if (p) diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index c713852fe93..afbc4e551d4 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -14455,6 +14455,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 (); return; } @@ -41432,6 +41437,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); @@ -41535,6 +41544,10 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok, return NULL_TREE; } + 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); @@ -41625,6 +41638,10 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok, return NULL_TREE; } + 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); @@ -43819,7 +43836,7 @@ cp_parser_omp_requires (cp_parser *parser, cp_token *pragma_tok) cp_parser_skip_to_pragma_eol (parser, pragma_tok); return false; } - if (p) + if (this_req == OMP_REQUIRES_DYNAMIC_ALLOCATORS) sorry_at (cloc, "%qs clause on % directive not " "supported yet", p); if (p) diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index cb166f956b7..c25531a4989 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -3668,7 +3668,7 @@ gfc_check_omp_requires (gfc_namespace *ns, int ref_omp_requires) if ((ref_omp_requires & OMP_REQ_REVERSE_OFFLOAD) && !(ns->omp_requires & OMP_REQ_REVERSE_OFFLOAD)) gfc_error ("Program unit at %L has OpenMP device constructs/routines " - "but does not set !$OMP REQUIRES REVERSE_OFFSET but other " + "but does not set !$OMP REQUIRES REVERSE_OFFLOAD but other " "program units do", &ns->proc_name->declared_at); if ((ref_omp_requires & OMP_REQ_UNIFIED_ADDRESS) && !(ns->omp_requires & OMP_REQ_UNIFIED_ADDRESS)) @@ -3855,7 +3855,8 @@ gfc_match_omp_requires (void) else goto error; - if (requires_clause & ~OMP_REQ_ATOMIC_MEM_ORDER_MASK) + /* Currently, everything except 'dynamic_allocators' is allowed. */ + if (requires_clause == 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)) diff --git a/gcc/fortran/parse.c b/gcc/fortran/parse.c index 1549f8e1635..4731bca2cf7 100644 --- a/gcc/fortran/parse.c +++ b/gcc/fortran/parse.c @@ -22,10 +22,12 @@ along with GCC; see the file COPYING3. If not see #include "system.h" #include "coretypes.h" #include "options.h" +#include "tree.h" #include "gfortran.h" #include #include "match.h" #include "parse.h" +#include "omp-general.h" /* Current statement label. Zero means no statement label. Because new_st can get wiped during statement matching, we have to keep it separate. */ @@ -6572,6 +6574,23 @@ done: gfc_current_ns = gfc_current_ns->sibling) gfc_check_omp_requires (gfc_current_ns, omp_requires); + if (omp_requires) + { + omp_requires_mask = (enum omp_requires) 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); + } + /* Do the parse tree dump. */ gfc_current_ns = flag_dump_fortran_original ? gfc_global_ns_list : NULL; diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index ba0937fba94..9cc7d2945fc 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -437,6 +437,24 @@ omp_finish_file (void) varpool_node::finalize_decl (vars_decl); varpool_node::finalize_decl (funcs_decl); + + if (flag_openmp && (omp_requires_mask & OMP_REQUIRES_TARGET_USED) != 0) + { + 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); + } } else { 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/gfortran.dg/gomp/requires-4.f90 b/gcc/testsuite/gfortran.dg/gomp/requires-4.f90 index b17aceb898b..c870a2840d3 100644 --- a/gcc/testsuite/gfortran.dg/gomp/requires-4.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/requires-4.f90 @@ -9,7 +9,7 @@ end module m subroutine foo !$omp target !$omp end target -! { dg-error "OpenMP device constructs/routines but does not set !.OMP REQUIRES REVERSE_OFFSET but other program units do" "" { target *-*-* } 9 } +! { dg-error "OpenMP device constructs/routines but does not set !.OMP REQUIRES REVERSE_OFFLOAD but other program units do" "" { target *-*-* } 9 } ! { dg-error "OpenMP device constructs/routines but does not set !.OMP REQUIRES UNIFIED_ADDRESS but other program units do" "" { target *-*-* } 9 } ! { dg-error "OpenMP device constructs/routines but does not set !.OMP REQUIRES UNIFIED_SHARED_MEMORY but other program units do" "" { target *-*-* } 9 } end diff --git a/gcc/testsuite/gfortran.dg/gomp/requires-8.f90 b/gcc/testsuite/gfortran.dg/gomp/requires-8.f90 index 3c32ae9860e..3819b0c28cc 100644 --- a/gcc/testsuite/gfortran.dg/gomp/requires-8.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/requires-8.f90 @@ -13,7 +13,7 @@ contains end subroutine foo end module m -subroutine bar ! { dg-error "has OpenMP device constructs/routines but does not set !.OMP REQUIRES REVERSE_OFFSET but other program units do" } +subroutine bar ! { dg-error "has OpenMP device constructs/routines but does not set !.OMP REQUIRES REVERSE_OFFLOAD but other program units do" } !use m !$omp requires unified_shared_memory !$omp declare target diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 11a9308e3d2..d5a0b2c5ea7 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -301,6 +301,12 @@ enum gomp_map_kind #define GOMP_DEPEND_INOUT 3 #define GOMP_DEPEND_MUTEXINOUTSET 4 +/* 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 b19428af6d8..78210a88f15 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 62645ce9954..f54469fdd6b 100644 --- a/libgomp/libgomp-plugin.h +++ b/libgomp/libgomp-plugin.h @@ -122,6 +122,7 @@ extern int GOMP_OFFLOAD_get_type (void); extern int GOMP_OFFLOAD_get_num_devices (void); extern bool GOMP_OFFLOAD_init_device (int); extern bool GOMP_OFFLOAD_fini_device (int); +extern bool GOMP_OFFLOAD_supported_features (unsigned *); extern unsigned GOMP_OFFLOAD_version (void); extern int GOMP_OFFLOAD_load_image (int, unsigned, const void *, struct addr_pair **); diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 305cba3aa02..09f2ac67943 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1130,6 +1130,7 @@ struct gomp_device_descr __typeof (GOMP_OFFLOAD_get_num_devices) *get_num_devices_func; __typeof (GOMP_OFFLOAD_init_device) *init_device_func; __typeof (GOMP_OFFLOAD_fini_device) *fini_device_func; + __typeof (GOMP_OFFLOAD_supported_features) *supported_features_func; __typeof (GOMP_OFFLOAD_version) *version_func; __typeof (GOMP_OFFLOAD_load_image) *load_image_func; __typeof (GOMP_OFFLOAD_unload_image) *unload_image_func; diff --git a/libgomp/oacc-host.c b/libgomp/oacc-host.c index f3bbd2b9c61..94a7fac2a39 100644 --- a/libgomp/oacc-host.c +++ b/libgomp/oacc-host.c @@ -71,6 +71,12 @@ host_fini_device (int n __attribute__ ((unused))) return true; } +static bool +host_supported_features (unsigned int *n) +{ + return (*n == 0); +} + static unsigned host_version (void) { @@ -273,6 +279,7 @@ static struct gomp_device_descr host_dispatch = .get_num_devices_func = host_get_num_devices, .init_device_func = host_init_device, .fini_device_func = host_fini_device, + .supported_features_func = host_supported_features, .version_func = host_version, .load_image_func = host_load_image, .unload_image_func = host_unload_image, diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c index 47f0b6e25f8..718d78173fe 100644 --- a/libgomp/plugin/plugin-gcn.c +++ b/libgomp/plugin/plugin-gcn.c @@ -3991,4 +3991,12 @@ GOMP_OFFLOAD_openacc_destroy_thread_data (void *data) free (data); } +/* Indicate which GOMP_REQUIRES_* features are supported, currently none. */ + +bool +GOMP_OFFLOAD_supported_features (unsigned int *mask) +{ + return (*mask == 0); +} + /* }}} */ diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c index 681c344b9c2..4cc25fbe232 100644 --- a/libgomp/plugin/plugin-nvptx.c +++ b/libgomp/plugin/plugin-nvptx.c @@ -1236,6 +1236,14 @@ GOMP_OFFLOAD_fini_device (int n) return true; } +/* Indicate which GOMP_REQUIRES_* features are supported, currently none. */ + +bool +GOMP_OFFLOAD_supported_features (unsigned int *mask) +{ + return (*mask == 0); +} + /* Return the libgomp version number we're compatible with. There is no requirement for cross-version compatibility. */ diff --git a/libgomp/target.c b/libgomp/target.c index 4a4e1f80745..f06df7ba28d 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -31,6 +31,7 @@ #include "gomp-constants.h" #include #include +#include #include #ifdef HAVE_INTTYPES_H # include /* For PRIu64. */ @@ -79,6 +80,16 @@ static int num_devices; /* Number of GOMP_OFFLOAD_CAP_OPENMP_400 devices. */ static int num_devices_openmp; +/* 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. */ +static unsigned int gomp_requires_mask; + +/* Start/end of .gnu.gomp.requires section of program, defined in + crtoffloadbegin/end.o. */ +extern const unsigned int __requires_mask_table[]; +extern const unsigned int __requires_mask_table_end[]; + /* Similar to gomp_realloc, but release register_lock before gomp_fatal. */ static void * @@ -1961,6 +1972,20 @@ gomp_init_device (struct gomp_device_descr *devicep) gomp_fatal ("device initialization failed"); } + unsigned int features = gomp_requires_mask; + if (!devicep->supported_features_func (&features)) + { + char buf[64], *end = buf + sizeof (buf), *p = buf; + if (features & GOMP_REQUIRES_UNIFIED_ADDRESS) + p += snprintf (p, end - p, "unified_address"); + if (features & GOMP_REQUIRES_UNIFIED_SHARED_MEMORY) + p += snprintf (p, end - p, "%sunified_shared_memory", + (p == buf ? "" : ", ")); + if (features & GOMP_REQUIRES_REVERSE_OFFLOAD) + p += snprintf (p, end - p, "%sreverse_offload", (p == buf ? "" : ", ")); + gomp_error ("device does not support required features: %s", buf); + } + /* Load to device all images registered by the moment. */ for (i = 0; i < num_offload_images; i++) { @@ -3200,6 +3225,7 @@ gomp_load_plugin_for_device (struct gomp_device_descr *device, DLSYM (get_num_devices); DLSYM (init_device); DLSYM (fini_device); + DLSYM (supported_features); DLSYM (load_image); DLSYM (unload_image); DLSYM (alloc); @@ -3310,6 +3336,28 @@ gomp_target_init (void) if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_DISABLED) return; + gomp_requires_mask = 0; + const unsigned int *mask_ptr = __requires_mask_table; + bool error_emitted = false; + while (mask_ptr != __requires_mask_table_end) + { + if (gomp_requires_mask == 0) + gomp_requires_mask = *mask_ptr; + else if (gomp_requires_mask != *mask_ptr) + { + if (!error_emitted) + { + gomp_error ("requires-directive clause inconsistency between " + "compilation units detected"); + error_emitted = true; + } + /* This is inconsistent, but still merge to query for all features + later. */ + gomp_requires_mask |= *mask_ptr; + } + mask_ptr++; + } + cur = OFFLOAD_PLUGINS; if (*cur) do 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..b5a3c512d28 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-1.c @@ -0,0 +1,21 @@ +/* { 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 "device does not support required features" } */ +/* { dg-shouldfail "" } */ 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..8b9341523c6 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-2-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-2.c b/libgomp/testsuite/libgomp.c-c++-common/requires-2.c new file mode 100644 index 00000000000..6fb280baabd --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-2.c @@ -0,0 +1,20 @@ +/* { dg-additional-sources requires-2-aux.c } */ + +#pragma omp requires 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; +} + +/* { dg-output "libgomp: device does not support required features: reverse_offload" } */ +/* { dg-shouldfail "" } */ diff --git a/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp index d1678d0514e..f92418fa416 100644 --- a/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp +++ b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp @@ -233,6 +233,14 @@ GOMP_OFFLOAD_fini_device (int device) return true; } +/* Indicate which GOMP_REQUIRES_* features are supported, currently none. */ + +extern "C" bool +GOMP_OFFLOAD_supported_features (unsigned int *mask) +{ + return (*mask == 0); +} + static bool get_target_table (int device, int &num_funcs, int &num_vars, void **&table) {