[OpenMP,5.0] More implementation of the requires directive

Hi Jakub,
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?

Thanks,
Chung-Lin

2021-01-13  Chung-Lin Tang  <cltang@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.

On Wed, Jan 13, 2021 at 11:07:44PM +0800, Chung-Lin Tang wrote:
> 2021-01-13  Chung-Lin Tang  <cltang@codesourcery.com>

...
Looks mostly ok, with some nits.

> 	* parse.c ("tree.h"): Add include.
> 	("omp-general.h"): Likewise.

I think the usual way is to write:
	* parse.c: Include "tree.h" and "omp-general.h".
 	(gfc_parse_file): Add code to merge omp_requires to omp_requires_mask.

Something I miss in the patch is that for the device API calls (I'd bother
only with direct calls) we should probably set OMP_REQUIRES_TARGET_USED
too, so perhaps do that during gimplification if flag_openmp and
in gimplify_call_expr there is fndecl and DECL_NAME of it is non-NULL and
starts with "omp_" it looks at DECL_ASSEMBLER_NAME and compares that to a
selected list of device APIs.

Also, would it be possible to diagnose .gnu.gomp_requires mismatches also
at link time through the linker plugin/mkoffload?
At least if we have LTO offloading bytecode in and the plugin is doing
something...

> +      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);

This probably needs to sorry if the target doesn't support named sections.
We probably don't support LTO in that case either though.

Also, the diagnostic of the mismatches on the library side should print
details, say that libfoobar is #pragma omp requires unified_shared_memory
while libbar is not.

	Jakub

Hi!

On 2021-01-13T23:07:44+0800, Chung-Lin Tang <cltang@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  <cltang@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

On 13/01/2021 15:07, Chung-Lin Tang wrote:
> 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.

After the Unified Shared Memory patches are committed, this patch will 
need to be altered as attached.

I'll commit my patch to OG11 shortly.

Andrew
libgomp, nvptx: report USM supported

libgomp/ChangeLog:

	* plugin/plugin-nvptx.c (GOMP_OFFLOAD_supported_features): Allow
	GOMP_REQUIRES_UNIFIED_ADDRESS and GOMP_REQUIRES_UNIFIED_SHARED_MEMORY.

diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index dd490b2ae2a..e77c6a87930 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -1260,11 +1260,14 @@ GOMP_OFFLOAD_fini_device (int n)
   return true;
 }
 
-/* Indicate which GOMP_REQUIRES_* features are supported, currently none.  */
+/* Indicate which GOMP_REQUIRES_* features are supported.  */
 
 bool
 GOMP_OFFLOAD_supported_features (unsigned int *mask)
 {
+  *mask &= ~(GOMP_REQUIRES_UNIFIED_ADDRESS
+             | GOMP_REQUIRES_UNIFIED_SHARED_MEMORY);
+
   return (*mask == 0);
 }