Hi!
On 2021-01-13T23:07:44+0800, Chung-Lin Tang <clt...@codesourcery.com> 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 <clt...@codesourcery.com>
>
> 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 (<stdio.h>): 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 %<requires%> 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 %<requires%> 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 <setjmp.h>
> #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 <limits.h>
> #include <stdbool.h>
> +#include <stdio.h>
> #include <stdlib.h>
> #ifdef HAVE_INTTYPES_H
> # include <inttypes.h> /* 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
>From ff77b4a0db75bc82a5519e31a882f9a25a02cd56 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Wed, 3 Mar 2021 22:37:58 +0100
Subject: [PATCH 1/2] [WIP] OpenMP 5.0: requires directive: workaround to fix
libgomp IntelMIC plugin build
Fix-up for og10 commit c2e4a17adc0989f216c7fc3f93f150c66adba23a "OpenMP 5.0:
requires directive".
The GCC offloading target configurations don't build/use
'crtoffloadbegin.o'/'crtoffloadtable.o'/'crtoffloadend.o'
('libgcc/offloadstuff.c'), but the libgomp IntelMIC plugin still does link
against libgomp, and the latter unconditionally refers to
'__requires_mask_table', '__requires_mask_table_end':
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 have not researched what a proper fix would look like.
libgomp/
* target.c (__requires_mask_table, __requires_mask_table_end): Add
'__attribute__((weak))'.
---
libgomp/ChangeLog.omp | 5 +++++
libgomp/target.c | 2 ++
2 files changed, 7 insertions(+)
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 0e3fd122f850..03ca74c8f3d5 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,8 @@
+2021-03-25 Thomas Schwinge <tho...@codesourcery.com>
+
+ * target.c (__requires_mask_table, __requires_mask_table_end): Add
+ '__attribute__((weak))'.
+
2021-03-01 Kwok Cheung Yeung <k...@codesourcery.com>
* testsuite/libgomp.c-c++-common/collapse-4.c: New.
diff --git a/libgomp/target.c b/libgomp/target.c
index 699dc72cb722..9c7582635aa3 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -104,7 +104,9 @@ static unsigned int gomp_requires_mask;
/* 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. */
--
2.30.2
>From 4ef4921cb10693c59b488002179db131683af8bc Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Wed, 3 Mar 2021 22:51:01 +0100
Subject: [PATCH 2/2] OpenMP 5.0: requires directive: adjust libgomp HSA plugin
Fix-up for og10 commit c2e4a17adc0989f216c7fc3f93f150c66adba23a "OpenMP 5.0:
requires directive".
libgomp: while loading libgomp-plugin-hsa.so.1: [...]/libgomp-plugin-hsa.so.1: undefined symbol: GOMP_OFFLOAD_supported_features
libgomp/
* plugin/plugin-hsa.c (GOMP_OFFLOAD_supported_features): New
function.
---
libgomp/ChangeLog.omp | 3 +++
libgomp/plugin/plugin-hsa.c | 8 ++++++++
2 files changed, 11 insertions(+)
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 03ca74c8f3d5..19f48dc61202 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,5 +1,8 @@
2021-03-25 Thomas Schwinge <tho...@codesourcery.com>
+ * plugin/plugin-hsa.c (GOMP_OFFLOAD_supported_features): New
+ function.
+
* target.c (__requires_mask_table, __requires_mask_table_end): Add
'__attribute__((weak))'.
diff --git a/libgomp/plugin/plugin-hsa.c b/libgomp/plugin/plugin-hsa.c
index abd3bc64163b..bddb690ca14f 100644
--- a/libgomp/plugin/plugin-hsa.c
+++ b/libgomp/plugin/plugin-hsa.c
@@ -1869,3 +1869,11 @@ GOMP_OFFLOAD_dev2dev (int ord, void *dst, const void *src, size_t n)
"it should never be called");
return false;
}
+
+/* Indicate which GOMP_REQUIRES_* features are supported, currently none. */
+
+bool
+GOMP_OFFLOAD_supported_features (unsigned int *mask)
+{
+ return (*mask == 0);
+}
--
2.30.2
