Hi! On 2021-01-13T23:07:44+0800, Chung-Lin Tang wrote: > this patch provides more implementation of the requires directive, basically: > > (1) The collection of the reverse_offload, unified_address, and unified_shared_memory > clauses into a .gnu.gomp_requires section > > (2) libgomp checking of consistency across the entire .gnu.gomp_requires section, > and querying into the offload plugin to see if the offload target supports the required > features (as of now, the setting is that none of those features are supported by any > of the plugins). > > We currently emit errors, but do not fatally cause exit of the program if those > are not met. We're still unsure if complete block-out of program execution is the right > thing for the user. This can be discussed later. > > Is this okay for trunk after stage1 re-opens? (As posted, per a quick check) this got pushed to devel/omp/gcc-10 branch in commit c2e4a17adc0989f216c7fc3f93f150c66adba23a "OpenMP 5.0: requires directive". Building the libgomp Intel MIC plugin fails: make[3]: Entering directory '[...]/build-gcc-offload-x86_64-intelmicemul-linux-gnu/x86_64-intelmicemul-linux-gnu/liboffloadmic/plugin' [...]/build-gcc-offload-x86_64-intelmicemul-linux-gnu/./gcc/xg++ [...] -loffloadmic_target -lcoi_device -lgomp -rdynamic ../ofldbegin.o offload_target_main.o ../ofldend.o -o offload_target_main ./../../libgomp/.libs/libgomp.so: undefined reference to `__requires_mask_table_end' ./../../libgomp/.libs/libgomp.so: undefined reference to `__requires_mask_table' collect2: error: ld returned 1 exit status Makefile:806: recipe for target 'offload_target_main' failed make[3]: *** [offload_target_main] Error 1 I've pushed "[WIP] OpenMP 5.0: requires directive: workaround to fix libgomp IntelMIC plugin build" to devel/omp/gcc-10 branch in commit ff77b4a0db75bc82a5519e31a882f9a25a02cd56, see attached. This seemed like a safe default, to get this un-stuck, but I suppose this will need further work. I haven't read up what this OpenMP functionality exactly is, and haven't thought about how it ought to be implemented -- but from a quick look, instead of libgomp directly referring to '__requires_mask_table', shouldn't this use some "dynamic indirection scheme" (like we have for the dynamic offloading code registering/loading function calls via constructors, synthesized by the 'mkoffload's?), so that it also works for shared objects ('*.so', etc.) containing OpenMP code? But maybe I just have no clue what I'm talking about, and this is not applicable here. ;-) 'make check-target-libgomp': libgomp: while loading libgomp-plugin-hsa.so.1: [...]/libgomp-plugin-hsa.so.1: undefined symbol: GOMP_OFFLOAD_supported_features I've pushed "OpenMP 5.0: requires directive: adjust libgomp HSA plugin" to devel/omp/gcc-10 branch in commit 4ef4921cb10693c59b488002179db131683af8bc, see attached. (The libgomp HSA plugin has been removed in master branch, so not applicable there.) Grüße Thomas > 2021-01-13 Chung-Lin Tang > > gcc/c/ > * c-parser.c (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): Adjust to only mention "not implemented yet" > for OMP_REQUIRES_DYNAMIC_ALLOCATORS. > > gcc/cp/ > * parser.c (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): Adjust to only mention "not implemented yet" > for OMP_REQUIRES_DYNAMIC_ALLOCATORS. > > gcc/fortran/ > * openmp.c (gfc_check_omp_requires): Fix REVERSE_OFFLOAD typo. > (gfc_match_omp_requires): Adjust to only mention "not implemented yet" > for OMP_REQUIRES_DYNAMIC_ALLOCATORS. > * parse.c ("tree.h"): Add include. > ("omp-general.h"): Likewise. > (gfc_parse_file): Add code to merge omp_requires to omp_requires_mask. > > gcc/ > * omp-offload.c (omp_finish_file): Add code to reate OpenMP requires > mask variable in .gnu.gomp_requires section if needed. > > gcc/testsuite/ > * c-c++-common/gomp/requires-4.c: Remove prune of "not supported yet". > * gcc/testsuite/gfortran.dg/gomp/requires-4.f90: Fix REVERSE_OFFLOAD typo. > * gcc/testsuite/gfortran.dg/gomp/requires-8.f90: Likewise. > > include/ > * gomp-constants.h (GOMP_REQUIRES_UNIFIED_ADDRESS): New symbol. > (GOMP_REQUIRES_UNIFIED_SHARED_MEMORY): Likewise. > (GOMP_REQUIRES_REVERSE_OFFLOAD): Likewise. > > libgcc/ > * offloadstuff.c (__requires_mask_table): New symbol to mark start of > .gnu.gomp_requires section. > (__requires_mask_table_end): New symbol to mark end of > .gnu.gomp_requires section. > > libgomp/ > * libgomp-plugin.h (GOMP_OFFLOAD_supported_features): New declaration. > * libgomp.h (struct gomp_device_descr): New 'supported_features_func' > plugin hook field. > * oacc-host.c (host_supported_features): New host hook function. > (host_dispatch): Initialize 'supported_features_func' host hook. > * plugin/plugin-gcn.c (GOMP_OFFLOAD_supported_features): New function. > * plugin/plugin-nvptx.c (GOMP_OFFLOAD_supported_features): Likewise. > * target.c (): Add include of standard header. > (gomp_requires_mask): New static variable. > (__requires_mask_table): New declaration. > (__requires_mask_table_end): Likewise. > (gomp_load_plugin_for_device): Add loading of 'supported_features' hook. > (gomp_target_init): Add code to summarize .gnu._gomp_requires section > mask values, emit error if inconsistency found. > > * testsuite/libgomp.c-c++-common/requires-1.c: New test. > * testsuite/libgomp.c-c++-common/requires-1-aux.c: New file linked with > above test. > * testsuite/libgomp.c-c++-common/requires-2.c: New test. > * testsuite/libgomp.c-c++-common/requires-2-aux.c: New file linked with > above test. > > liboffloadmic/ > * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_supported_features): > New function. > 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) > { ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf