OpenMP 5.0: requires directive

This is a merge of:
https://gcc.gnu.org/pipermail/gcc-patches/2021-January/563393.html

This patch completes more of the reverse_offload, unified_address, and
unified_shared_memory clauses for the OpenMP 5.0 requires directive,
including runtime verification of the offload target.
(currently no offload devices actually support above features, only
warning messages are emitted)

This may possibly reverted/updated when a final patch is approved
for mainline.

2021-02-02  Chung-Lin Tang  <cltang@codesourcery.com>

gcc/c/ChangeLog:

	* 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/ChangeLog:

	* 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/ChangeLog:

	* 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/ChangeLog:

	* omp-offload.c (omp_finish_file): Add code to reate OpenMP requires
	mask variable in .gnu.gomp_requires section if needed.

gcc/testsuite/ChangeLog:

	* 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/ChangeLog:

	* gomp-constants.h (GOMP_REQUIRES_UNIFIED_ADDRESS): New symbol.
	(GOMP_REQUIRES_UNIFIED_SHARED_MEMORY): Likewise.
	(GOMP_REQUIRES_REVERSE_OFFLOAD): Likewise.

libgcc/ChangeLog:

	* 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/ChangeLog:

	* 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/ChangeLog:

	* 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 2679b3e6405e75e74f5c02a66a1e48a68b43098d..8b75e7aab6389858b47534adae8e1fe052971546 100644 (file)
--- 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
@@ -19625,6 +19631,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");
@@ -19767,6 +19777,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");
@@ -19853,6 +19867,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");
@@ -21440,7 +21458,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 3e33648180fdd67169b7a03e6d0deeccbc0b674b..f9a2092e0e195020b76550e22ba22637e0f7ae11 100644 (file)
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -14531,6 +14531,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;
            }
@@ -41623,6 +41628,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);
@@ -41726,6 +41735,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);
@@ -41816,6 +41829,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);
@@ -43971,7 +43988,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 dca2a782d7398fd4bb2f02678cc2386dbcc7ac23..c309513012d7ee252552510938e6e50cdb931d13 100644 (file)
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -3721,7 +3721,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))
@@ -3908,7 +3908,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 f45a6b71ccef4daf26224be7a285c9aba1a03441..8fc7678593a0f5ce2fe2dea29d36e2aa6efc8c14 100644 (file)
--- 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 543297e6b9e68e480faf53aa99169dba0dfc3ab4..821b60a392380c183669f9fa3a97603424a5ce93 100644 (file)
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -439,6 +439,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 88ba7746cf81698c3aeb3f9ed877448fd5ca43e0..8f45d83ea6e30c2d564ec14d6fd5f463575e1a51 100644 (file)
--- 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 b17aceb898b43dec9050155402395b657365e613..c870a2840d3ccc0ebb6dcc0b610097e00a47ee4b 100644 (file)
--- 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 3c32ae9860e57cb320d1967e8747edf374a77f0e..3819b0c28cc279e1c628eb03b8e4f67ba5023b8b 100644 (file)
--- 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 b11ed39fbff79a745ec690477bc7d88a96935662..d9b80ac97327c3aaf1f1a16b00f641f8a9fbf38e 100644 (file)
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -342,6 +342,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 b19428af6d879dcdeac48c09ea80cc47a206e937..78210a88f158ce908f1ecd54532c2fcd50983ddf 100644 (file)
--- 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 bff2193dd3a7c7a50c111b39a1b2a41c43a522f5..1b28f57b7b7a1e69ae867a9b8df4dd8b709ee9bb 100644 (file)
--- 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 4c476db09c8a511b79b036d459088dbc43effb7a..0c8085e350ca20283b2007df55cbf6174d09ca3a 100644 (file)
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1184,6 +1184,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 369abfffac9922ab5888f49ec94bebb44ee77ea3..1a0ff16a65ecc857e3e2325b40c80bf9e35526c4 100644 (file)
--- 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 002921cca3821cd64ac863a1d27b25daf6edcb50..1402e85db7497c1435c1996822d928656096512a 100644 (file)
--- a/libgomp/plugin/plugin-gcn.c
+++ b/libgomp/plugin/plugin-gcn.c
@@ -4068,4 +4068,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 82bf97948c6fc94079131f420f9277697bae2205..23573416a12e03c2ce40c6fac073586d85953e6a 100644 (file)
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -1234,6 +1234,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 de551da4f635d36e187940b6ff8d09a9e7a288df..af39d283608c1bb130e133afadda7c7b44a648ee 100644 (file)
--- 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.  */
@@ -96,6 +97,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 *
@@ -2384,6 +2395,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++)
     {
@@ -3664,6 +3689,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);
@@ -3774,6 +3800,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 (file)
index 0000000..8b93415
--- /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 (file)
index 0000000..02585ad
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-1.c
@@ -0,0 +1,20 @@
+/* { 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" } */

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 (file)
index 0000000..8b93415
--- /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 (file)
index 0000000..0302f89
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-2.c
@@ -0,0 +1,19 @@
+/* { 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" } */

diff --git a/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
index d1678d0514e91e775fff7fe361b861f706b53782..f92418fa4164b9fb68a8081982764d21d725f6ba 100644 (file)
--- 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)
 {