public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-12] OpenMP: Move omp requires checks to libgomp
@ 2022-07-07 6:56 Tobias Burnus
0 siblings, 0 replies; only message in thread
From: Tobias Burnus @ 2022-07-07 6:56 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:5e7baa3c8875018d94e4222bb26c42a72c7499cb
commit 5e7baa3c8875018d94e4222bb26c42a72c7499cb
Author: Tobias Burnus <tobias@codesourcery.com>
Date: Tue Jul 5 15:31:01 2022 +0200
OpenMP: Move omp requires checks to libgomp
This commit reverts OG12 commits
commit 6ab6303f61c812d6e2f05d44af5cc79815c57bb8
'OpenMP 5.0: requires directive'
and
commit 47e5ba646c20fdb64fa233dd570f155880cafd04
'[WIP] OpenMP 5.0: requires directive: workaround to fix libgomp IntelMIC plugin build'
And replaces those by the upstream version, commit r13-1458, i.e.
(cherry picked from commit 683f11843974f0bdf42f79cdcbb0c2b43c7b81b0)
It also updates
* libgomp/plugin/plugin-gcn.c (GOMP_OFFLOAD_get_num_devices)
to permit GOMP_REQUIRES_UNIFIED_ADDRESS | GOMP_REQUIRES_UNIFIED_SHARED_MEMORY,
moved from the now no-longer existing function GOMP_OFFLOAD_supported_features;
the flags were added in OG12 commit 12d14a9a255c1cc10e4506935327aabd9766967d,
'amdgcn: libgomp plugin USM implementation'.
And, likewise, it updates
* libgomp/plugin/plugin-nvptx.c (GOMP_OFFLOAD_get_num_devices)
from OG12 commit f3fd38e31fafdc7f6b6a721b53bdc595c4ec1e09,
'libgomp, nvptx: report USM'.
* * *
Handle reverse_offload, unified_address, and unified_shared_memory
requirements in libgomp by saving them alongside the offload table.
When the device lto1 runs, it extracts the data for mkoffload. The
latter than passes the value on to GOMP_offload_register_ver.
lto1 (either the host one, with -flto [+ ENABLE_OFFLOADING], or in the
offload-device lto1) also does the the consistency check is done,
erroring out when the 'omp requires' clause use is inconsistent.
For all in-principle supported devices, if a requirement cannot be fulfilled,
the device is excluded from the (supported) devices list. Currently, none of
those requirements are marked as supported for any of the non-host devices.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_target_data, c_parser_omp_target_update,
c_parser_omp_target_enter_data, c_parser_omp_target_exit_data): Set
OMP_REQUIRES_TARGET_USED.
(c_parser_omp_requires): Remove sorry.
gcc/ChangeLog:
* config/gcn/mkoffload.cc (process_asm): Write '#include <stdint.h>'.
(process_obj): Pass omp_requires_mask to GOMP_offload_register_ver.
(main): Ask lto1 to obtain omp_requires_mask and pass it on.
* config/nvptx/mkoffload.cc (process, main): Likewise.
* lto-cgraph.cc (omp_requires_to_name): New.
(input_offload_tables): Save omp_requires_mask.
(output_offload_tables): Read it, check for consistency,
save value for mkoffload.
* omp-low.cc (lower_omp_target): Force output_offloadtables
call for OMP_REQUIRES_TARGET_USED.
gcc/cp/ChangeLog:
* parser.cc (cp_parser_omp_target_data,
cp_parser_omp_target_enter_data, cp_parser_omp_target_exit_data,
cp_parser_omp_target_update): Set OMP_REQUIRES_TARGET_USED.
(cp_parser_omp_requires): Remove sorry.
gcc/fortran/ChangeLog:
* openmp.cc (gfc_match_omp_requires): Remove sorry.
* parse.cc (decode_omp_directive): Don't regard 'declare target'
as target usage for 'omp requires'; add more flags to
omp_requires_mask.
include/ChangeLog:
* gomp-constants.h (GOMP_VERSION): Bump to 2.
(GOMP_REQUIRES_UNIFIED_ADDRESS, GOMP_REQUIRES_UNIFIED_SHARED_MEMORY,
GOMP_REQUIRES_REVERSE_OFFLOAD, GOMP_REQUIRES_TARGET_USED):
New defines.
libgomp/ChangeLog:
* libgomp-plugin.h (GOMP_OFFLOAD_get_num_devices): Add
omp_requires_mask arg.
* plugin/plugin-gcn.c (GOMP_OFFLOAD_get_num_devices): Likewise;
return -1 when device available but omp_requires_mask != 0.
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_get_num_devices): Likewise.
* oacc-host.c (host_get_num_devices, host_openacc_get_property):
Update call.
* oacc-init.c (resolve_device, acc_init_1, acc_shutdown_1,
goacc_attach_host_thread_to_device, acc_get_num_devices,
acc_set_device_num, get_property_any): Likewise.
* target.c (omp_requires_mask): New global var.
(gomp_requires_to_name): New.
(GOMP_offload_register_ver): Handle passed omp_requires_mask.
(gomp_target_init): Handle omp_requires_mask.
* libgomp.texi (OpenMP 5.0): Update requires impl. status.
(OpenMP 5.1): Add a missed item.
(OpenMP 5.2): Mark linear-clause change as supported in C/C++.
* testsuite/libgomp.c-c++-common/requires-1-aux.c: New test.
* testsuite/libgomp.c-c++-common/requires-1.c: New test.
* testsuite/libgomp.c-c++-common/requires-2-aux.c: New test.
* testsuite/libgomp.c-c++-common/requires-2.c: New test.
* testsuite/libgomp.c-c++-common/requires-3-aux.c: New test.
* testsuite/libgomp.c-c++-common/requires-3.c: New test.
* testsuite/libgomp.c-c++-common/requires-4-aux.c: New test.
* testsuite/libgomp.c-c++-common/requires-4.c: New test.
* testsuite/libgomp.c-c++-common/requires-5-aux.c: New test.
* testsuite/libgomp.c-c++-common/requires-5.c: New test.
* testsuite/libgomp.c-c++-common/requires-6.c: New test.
* testsuite/libgomp.c-c++-common/requires-7-aux.c: New test.
* testsuite/libgomp.c-c++-common/requires-7.c: New test.
* testsuite/libgomp.fortran/requires-1-aux.f90: New test.
* testsuite/libgomp.fortran/requires-1.f90: New test.
liboffloadmic/ChangeLog:
* plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_num_devices):
Return -1 when device available but omp_requires_mask != 0.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/requires-4.c: Update dg-*.
* c-c++-common/gomp/reverse-offload-1.c: Likewise.
* c-c++-common/gomp/target-device-ancestor-2.c: Likewise.
* c-c++-common/gomp/target-device-ancestor-3.c: Likewise.
* c-c++-common/gomp/target-device-ancestor-4.c: Likewise.
* c-c++-common/gomp/target-device-ancestor-5.c: Likewise.
* gfortran.dg/gomp/target-device-ancestor-3.f90: Likewise.
* gfortran.dg/gomp/target-device-ancestor-4.f90: Likewise.
* gfortran.dg/gomp/target-device-ancestor-5.f90: Likewise.
* gfortran.dg/gomp/target-device-ancestor-2.f90: Likewise. Move
post-FE checks to ...
* gfortran.dg/gomp/target-device-ancestor-2a.f90: ... this new file.
* gfortran.dg/gomp/requires-8.f90: Update as we don't regard
'declare target' for the 'requires' usage requirement.
Co-authored-by: Chung-Lin Tang <cltang@codesourcery.com>
Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
Diff:
---
gcc/ChangeLog.omp | 26 +++++
gcc/c/ChangeLog.omp | 21 ++++
gcc/c/c-parser.cc | 10 +-
gcc/config/gcn/mkoffload.cc | 27 ++++-
gcc/config/nvptx/mkoffload.cc | 29 ++++-
gcc/cp/ChangeLog.omp | 21 ++++
gcc/cp/parser.cc | 9 +-
gcc/fortran/ChangeLog.omp | 12 ++
gcc/fortran/openmp.cc | 2 +-
gcc/fortran/parse.cc | 40 +++----
gcc/lto-cgraph.cc | 117 ++++++++++++++++++-
gcc/omp-low.cc | 5 +
gcc/omp-offload.cc | 18 ---
gcc/testsuite/ChangeLog.omp | 22 ++++
.../c-c++-common/gomp/reverse-offload-1.c | 2 +-
.../c-c++-common/gomp/target-device-ancestor-2.c | 8 +-
.../c-c++-common/gomp/target-device-ancestor-5.c | 2 +-
gcc/testsuite/gfortran.dg/gomp/requires-8.f90 | 14 ++-
.../gfortran.dg/gomp/target-device-ancestor-2a.f90 | 80 +++++++++++++
.../gfortran.dg/gomp/target-device-ancestor-5.f90 | 8 +-
include/ChangeLog.omp | 10 ++
include/gomp-constants.h | 3 +-
libgcc/ChangeLog.omp | 10 ++
libgcc/offloadstuff.c | 9 --
libgomp/ChangeLog.omp | 65 +++++++++++
libgomp/libgomp-plugin.h | 3 +-
libgomp/libgomp.h | 1 -
libgomp/libgomp.texi | 7 +-
libgomp/oacc-host.c | 11 +-
libgomp/oacc-init.c | 16 +--
libgomp/plugin/plugin-gcn.c | 19 ++--
libgomp/plugin/plugin-nvptx.c | 24 ++--
libgomp/target.c | 124 ++++++++++++---------
.../libgomp.c-c++-common/requires-1-aux.c | 2 +-
.../testsuite/libgomp.c-c++-common/requires-1.c | 9 +-
.../libgomp.c-c++-common/requires-2-aux.c | 2 -
.../testsuite/libgomp.c-c++-common/requires-2.c | 10 +-
.../libgomp.c-c++-common/requires-3-aux.c | 11 ++
.../testsuite/libgomp.c-c++-common/requires-3.c | 24 ++++
.../libgomp.c-c++-common/requires-4-aux.c | 13 +++
.../testsuite/libgomp.c-c++-common/requires-4.c | 23 ++++
.../libgomp.c-c++-common/requires-5-aux.c | 11 ++
.../testsuite/libgomp.c-c++-common/requires-5.c | 21 ++++
.../testsuite/libgomp.c-c++-common/requires-6.c | 17 +++
.../libgomp.c-c++-common/requires-7-aux.c | 11 ++
.../testsuite/libgomp.c-c++-common/requires-7.c | 24 ++++
.../testsuite/libgomp.fortran/requires-1-aux.f90 | 14 +++
libgomp/testsuite/libgomp.fortran/requires-1.f90 | 26 +++++
liboffloadmic/ChangeLog.omp | 18 +++
liboffloadmic/plugin/libgomp-plugin-intelmic.cpp | 14 +--
50 files changed, 826 insertions(+), 199 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 8da5857cc63..a7bdb1c41fe 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,29 @@
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-07-04 Tobias Burnus <tobias@codesourcery.com>
+ Chung-Lin Tang <cltang@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * config/gcn/mkoffload.cc (process_asm): Write '#include <stdint.h>'.
+ (process_obj): Pass omp_requires_mask to GOMP_offload_register_ver.
+ (main): Ask lto1 to obtain omp_requires_mask and pass it on.
+ * config/nvptx/mkoffload.cc (process, main): Likewise.
+ * lto-cgraph.cc (omp_requires_to_name): New.
+ (input_offload_tables): Save omp_requires_mask.
+ (output_offload_tables): Read it, check for consistency,
+ save value for mkoffload.
+ * omp-low.cc (lower_omp_target): Force output_offloadtables
+ call for OMP_REQUIRES_TARGET_USED.
+
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Reverted:
+ 2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * omp-offload.cc (omp_finish_file): Add code to create OpenMP requires
+ mask variable in .gnu.gomp_requires section if needed.
+
2022-07-05 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline:
diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp
index cf45947a0a4..31499fcc9c7 100644
--- a/gcc/c/ChangeLog.omp
+++ b/gcc/c/ChangeLog.omp
@@ -1,3 +1,24 @@
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-07-04 Tobias Burnus <tobias@codesourcery.com>
+ Chung-Lin Tang <cltang@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-parser.cc (c_parser_omp_target_data, c_parser_omp_target_update,
+ c_parser_omp_target_enter_data, c_parser_omp_target_exit_data): Set
+ OMP_REQUIRES_TARGET_USED.
+ (c_parser_omp_requires): Remove sorry.
+
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Reverted:
+ 2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * c-parser.cc (c_parser_declaration_or_fndef): Set
+ OMP_REQUIRES_TARGET_USED in omp_requires_mask if function has
+ "omp declare target" attribute.
+
2022-07-05 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline:
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 2ef391dbbb0..cec19865ea4 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -2503,12 +2503,6 @@ 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
@@ -21260,6 +21254,10 @@ c_parser_omp_target_update (location_t loc, c_parser *parser,
return false;
}
+ if (flag_openmp)
+ omp_requires_mask
+ = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
tree stmt = make_node (OMP_TARGET_UPDATE);
TREE_TYPE (stmt) = void_type_node;
OMP_TARGET_UPDATE_CLAUSES (stmt) = clauses;
diff --git a/gcc/config/gcn/mkoffload.cc b/gcc/config/gcn/mkoffload.cc
index c430c4e8cee..5741d0a917b 100644
--- a/gcc/config/gcn/mkoffload.cc
+++ b/gcc/config/gcn/mkoffload.cc
@@ -621,6 +621,7 @@ process_asm (FILE *in, FILE *out, FILE *cfile)
struct regcount *regcounts = XOBFINISH (®counts_os, struct regcount *);
fprintf (cfile, "#include <stdlib.h>\n");
+ fprintf (cfile, "#include <stdint.h>\n");
fprintf (cfile, "#include <stdbool.h>\n\n");
fprintf (cfile, "#include <stdio.h>\n\n");
@@ -703,7 +704,7 @@ process_asm (FILE *in, FILE *out, FILE *cfile)
/* Embed an object file into a C source file. */
static void
-process_obj (FILE *in, FILE *cfile)
+process_obj (FILE *in, FILE *cfile, uint32_t omp_requires)
{
size_t len = 0;
const char *input = read_file (in, &len);
@@ -731,16 +732,18 @@ process_obj (FILE *in, FILE *cfile)
fprintf (cfile,
"static const struct gcn_image_desc {\n"
+ " uintptr_t omp_requires_mask;\n"
" const struct gcn_image *gcn_image;\n"
" unsigned kernel_count;\n"
" const struct hsa_kernel_description *kernel_infos;\n"
" unsigned global_variable_count;\n"
"} target_data = {\n"
+ " %d,\n"
" &gcn_image,\n"
" sizeof (gcn_kernels) / sizeof (gcn_kernels[0]),\n"
" gcn_kernels,\n"
" gcn_num_vars\n"
- "};\n\n");
+ "};\n\n", omp_requires);
fprintf (cfile,
"#ifdef __cplusplus\n"
@@ -1119,9 +1122,27 @@ main (int argc, char **argv)
unsetenv ("COMPILER_PATH");
unsetenv ("LIBRARY_PATH");
+ char *omp_requires_file;
+ if (save_temps)
+ omp_requires_file = concat (dumppfx, ".mkoffload.omp_requires", NULL);
+ else
+ omp_requires_file = make_temp_file (".mkoffload.omp_requires");
+
/* Run the compiler pass. */
+ xputenv (concat ("GCC_OFFLOAD_OMP_REQUIRES_FILE=", omp_requires_file, NULL));
fork_execute (cc_argv[0], CONST_CAST (char **, cc_argv), true, ".gcc_args");
obstack_free (&cc_argv_obstack, NULL);
+ unsetenv("GCC_OFFLOAD_OMP_REQUIRES_FILE");
+
+ in = fopen (omp_requires_file, "rb");
+ if (!in)
+ fatal_error (input_location, "cannot open omp_requires file %qs",
+ omp_requires_file);
+ uint32_t omp_requires;
+ if (fread (&omp_requires, sizeof (omp_requires), 1, in) != 1)
+ fatal_error (input_location, "cannot read omp_requires file %qs",
+ omp_requires_file);
+ fclose (in);
in = fopen (gcn_s1_name, "r");
if (!in)
@@ -1144,7 +1165,7 @@ main (int argc, char **argv)
if (!in)
fatal_error (input_location, "cannot open intermediate gcn obj file");
- process_obj (in, cfile);
+ process_obj (in, cfile, omp_requires);
fclose (in);
diff --git a/gcc/config/nvptx/mkoffload.cc b/gcc/config/nvptx/mkoffload.cc
index b28c1a32292..d8c81eb0547 100644
--- a/gcc/config/nvptx/mkoffload.cc
+++ b/gcc/config/nvptx/mkoffload.cc
@@ -231,7 +231,7 @@ access_check (const char *name, int mode)
}
static void
-process (FILE *in, FILE *out)
+process (FILE *in, FILE *out, uint32_t omp_requires)
{
size_t len = 0;
const char *input = read_file (in, &len);
@@ -240,6 +240,8 @@ process (FILE *in, FILE *out)
unsigned obj_count = 0;
unsigned ix;
+ fprintf (out, "#include <stdint.h>\n\n");
+
/* Dump out char arrays for each PTX object file. These are
terminated by a NUL. */
for (size_t i = 0; i != len;)
@@ -309,6 +311,7 @@ process (FILE *in, FILE *out)
fprintf (out,
"static const struct nvptx_tdata {\n"
+ " uintptr_t omp_requires_mask;\n"
" const struct ptx_obj *ptx_objs;\n"
" unsigned ptx_num;\n"
" const char *const *var_names;\n"
@@ -316,12 +319,12 @@ process (FILE *in, FILE *out)
" const struct nvptx_fn *fn_names;\n"
" unsigned fn_num;\n"
"} target_data = {\n"
- " ptx_objs, sizeof (ptx_objs) / sizeof (ptx_objs[0]),\n"
+ " %d, ptx_objs, sizeof (ptx_objs) / sizeof (ptx_objs[0]),\n"
" var_mappings,"
" sizeof (var_mappings) / sizeof (var_mappings[0]),\n"
" func_mappings,"
" sizeof (func_mappings) / sizeof (func_mappings[0])\n"
- "};\n\n");
+ "};\n\n", omp_requires);
fprintf (out, "#ifdef __cplusplus\n"
"extern \"C\" {\n"
@@ -583,19 +586,37 @@ main (int argc, char **argv)
unsetenv ("COMPILER_PATH");
unsetenv ("LIBRARY_PATH");
+ char *omp_requires_file;
+ if (save_temps)
+ omp_requires_file = concat (dumppfx, ".mkoffload.omp_requires", NULL);
+ else
+ omp_requires_file = make_temp_file (".mkoffload.omp_requires");
+
+ xputenv (concat ("GCC_OFFLOAD_OMP_REQUIRES_FILE=", omp_requires_file, NULL));
fork_execute (new_argv[0], CONST_CAST (char **, new_argv), true,
".gcc_args");
obstack_free (&argv_obstack, NULL);
+ unsetenv("GCC_OFFLOAD_OMP_REQUIRES_FILE");
xputenv (concat ("GCC_EXEC_PREFIX=", execpath, NULL));
xputenv (concat ("COMPILER_PATH=", cpath, NULL));
xputenv (concat ("LIBRARY_PATH=", lpath, NULL));
+ in = fopen (omp_requires_file, "rb");
+ if (!in)
+ fatal_error (input_location, "cannot open omp_requires file %qs",
+ omp_requires_file);
+ uint32_t omp_requires;
+ if (fread (&omp_requires, sizeof (omp_requires), 1, in) != 1)
+ fatal_error (input_location, "cannot read omp_requires file %qs",
+ omp_requires_file);
+ fclose (in);
+
in = fopen (ptx_name, "r");
if (!in)
fatal_error (input_location, "cannot open intermediate ptx file");
- process (in, out);
+ process (in, out, omp_requires);
fclose (in);
}
diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp
index 0f29357f562..b856f4ae0d4 100644
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,24 @@
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-07-04 Tobias Burnus <tobias@codesourcery.com>
+ Chung-Lin Tang <cltang@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * parser.cc (cp_parser_omp_target_data,
+ cp_parser_omp_target_enter_data, cp_parser_omp_target_exit_data,
+ cp_parser_omp_target_update): Set OMP_REQUIRES_TARGET_USED.
+ (cp_parser_omp_requires): Remove sorry.
+
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Reverted:
+ 2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * parser.cc (cp_parser_simple_declaration): Set
+ OMP_REQUIRES_TARGET_USED in omp_requires_mask if function has
+ "omp declare target" attribute.
+
2022-07-05 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline:
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index e685d5f7c7b..ebdace5f6e9 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -15376,11 +15376,6 @@ 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 ();
cp_finalize_omp_declare_simd (parser, &odsd);
return;
@@ -44712,6 +44707,10 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok,
return true;
}
+ if (flag_openmp)
+ omp_requires_mask
+ = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
tree stmt = make_node (OMP_TARGET_UPDATE);
TREE_TYPE (stmt) = void_type_node;
OMP_TARGET_UPDATE_CLAUSES (stmt) = clauses;
diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index b4c6f6ec820..5e1b91afa82 100644
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,15 @@
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-07-04 Tobias Burnus <tobias@codesourcery.com>
+ Chung-Lin Tang <cltang@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * openmp.cc (gfc_match_omp_requires): Remove sorry.
+ * parse.cc (decode_omp_directive): Don't regard 'declare target'
+ as target usage for 'omp requires'; add more flags to
+ omp_requires_mask.
+
2022-07-05 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline:
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 2f17b686365..88f1ae63f15 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -6132,7 +6132,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_OFFLOAD but other "
+ "but does not set !$OMP REQUIRES REVERSE_OFFSET 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))
diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc
index 25fab8178c5..c0e84956f05 100644
--- a/gcc/fortran/parse.cc
+++ b/gcc/fortran/parse.cc
@@ -1183,7 +1183,8 @@ decode_omp_directive (void)
}
switch (ret)
{
- case ST_OMP_DECLARE_TARGET:
+ /* Set omp_target_seen; exclude ST_OMP_DECLARE_TARGET.
+ FIXME: Get clarification, cf. OpenMP Spec Issue #3240. */
case ST_OMP_TARGET:
case ST_OMP_TARGET_DATA:
case ST_OMP_TARGET_ENTER_DATA:
@@ -6980,34 +6981,20 @@ done:
/* Fixup for external procedures and resolve 'omp requires'. */
int omp_requires;
+ bool omp_target_seen;
omp_requires = 0;
+ omp_target_seen = false;
for (gfc_current_ns = gfc_global_ns_list; gfc_current_ns;
gfc_current_ns = gfc_current_ns->sibling)
{
omp_requires |= gfc_current_ns->omp_requires;
+ omp_target_seen |= gfc_current_ns->omp_target_seen;
gfc_check_externals (gfc_current_ns);
}
for (gfc_current_ns = gfc_global_ns_list; gfc_current_ns;
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);
- }
-
/* Populate omp_requires_mask (needed for resolving OpenMP
metadirectives and declare variant). */
switch (omp_requires & OMP_REQ_ATOMIC_MEM_ORDER_MASK)
@@ -7026,9 +7013,22 @@ done:
break;
}
- if (omp_requires & OMP_REQ_DYNAMIC_ALLOCATORS)
+ if (omp_target_seen)
+ omp_requires_mask = (enum omp_requires) (omp_requires_mask
+ | 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_DYNAMIC_ALLOCATORS);
+ = (enum omp_requires) (omp_requires_mask
+ | OMP_REQUIRES_UNIFIED_SHARED_MEMORY);
+ if (omp_requires & OMP_REQ_DYNAMIC_ALLOCATORS)
+ omp_requires_mask = (enum omp_requires) (omp_requires_mask
+ | OMP_REQUIRES_DYNAMIC_ALLOCATORS);
/* Do the parse tree dump. */
gfc_current_ns = flag_dump_fortran_original ? gfc_global_ns_list : NULL;
diff --git a/gcc/lto-cgraph.cc b/gcc/lto-cgraph.cc
index 237743ef0ba..48629651e31 100644
--- a/gcc/lto-cgraph.cc
+++ b/gcc/lto-cgraph.cc
@@ -37,6 +37,7 @@ along with GCC; see the file COPYING3. If not see
#include "pass_manager.h"
#include "ipa-utils.h"
#include "omp-offload.h"
+#include "omp-general.h"
#include "stringpool.h"
#include "attribs.h"
#include "alloc-pool.h"
@@ -1068,7 +1069,10 @@ read_string (class lto_input_block *ib)
void
output_offload_tables (void)
{
- if (vec_safe_is_empty (offload_funcs) && vec_safe_is_empty (offload_vars))
+ bool output_requires = (flag_openmp
+ && (omp_requires_mask & OMP_REQUIRES_TARGET_USED) != 0);
+ if (vec_safe_is_empty (offload_funcs) && vec_safe_is_empty (offload_vars)
+ && !output_requires)
return;
struct lto_simple_output_block *ob
@@ -1098,6 +1102,19 @@ output_offload_tables (void)
(*offload_vars)[i]);
}
+ if (output_requires)
+ {
+ HOST_WIDE_INT val = ((HOST_WIDE_INT) omp_requires_mask
+ & (OMP_REQUIRES_UNIFIED_ADDRESS
+ | OMP_REQUIRES_UNIFIED_SHARED_MEMORY
+ | OMP_REQUIRES_REVERSE_OFFLOAD
+ | OMP_REQUIRES_TARGET_USED));
+ /* (Mis)use LTO_symtab_edge for this variable. */
+ streamer_write_enum (ob->main_stream, LTO_symtab_tags,
+ LTO_symtab_last_tag, LTO_symtab_edge);
+ streamer_write_hwi_stream (ob->main_stream, val);
+ }
+
streamer_write_uhwi_stream (ob->main_stream, 0);
lto_destroy_simple_output_block (ob);
@@ -1764,6 +1781,20 @@ input_symtab (void)
}
}
+static void
+omp_requires_to_name (char *buf, size_t size, HOST_WIDE_INT requires_mask)
+{
+ char *end = buf + size, *p = buf;
+ if (requires_mask & GOMP_REQUIRES_UNIFIED_ADDRESS)
+ p += snprintf (p, end - p, "unified_address");
+ if (requires_mask & GOMP_REQUIRES_UNIFIED_SHARED_MEMORY)
+ p += snprintf (p, end - p, "%sunified_shared_memory",
+ (p == buf ? "" : ", "));
+ if (requires_mask & GOMP_REQUIRES_REVERSE_OFFLOAD)
+ p += snprintf (p, end - p, "%sreverse_offload",
+ (p == buf ? "" : ", "));
+}
+
/* Input function/variable tables that will allow libgomp to look up offload
target code, and store them into OFFLOAD_FUNCS and OFFLOAD_VARS. */
@@ -1773,6 +1804,10 @@ input_offload_tables (bool do_force_output)
struct lto_file_decl_data **file_data_vec = lto_get_file_decl_data ();
struct lto_file_decl_data *file_data;
unsigned int j = 0;
+ const char *requires_fn = NULL;
+ tree requires_decl = NULL_TREE;
+
+ omp_requires_mask = (omp_requires) 0;
while ((file_data = file_data_vec[j++]))
{
@@ -1784,6 +1819,7 @@ input_offload_tables (bool do_force_output)
if (!ib)
continue;
+ tree tmp_decl = NULL_TREE;
enum LTO_symtab_tags tag
= streamer_read_enum (ib, LTO_symtab_tags, LTO_symtab_last_tag);
while (tag)
@@ -1799,6 +1835,7 @@ input_offload_tables (bool do_force_output)
LTO mode. */
if (do_force_output)
cgraph_node::get (fn_decl)->mark_force_output ();
+ tmp_decl = fn_decl;
}
else if (tag == LTO_symtab_variable)
{
@@ -1810,6 +1847,72 @@ input_offload_tables (bool do_force_output)
may be no refs to var_decl in offload LTO mode. */
if (do_force_output)
varpool_node::get (var_decl)->force_output = 1;
+ tmp_decl = var_decl;
+ }
+ else if (tag == LTO_symtab_edge)
+ {
+ static bool error_emitted = false;
+ HOST_WIDE_INT val = streamer_read_hwi (ib);
+
+ if (omp_requires_mask == 0)
+ {
+ omp_requires_mask = (omp_requires) val;
+ requires_decl = tmp_decl;
+ requires_fn = file_data->file_name;
+ }
+ else if (omp_requires_mask != val && !error_emitted)
+ {
+ const char *fn1 = requires_fn;
+ if (requires_decl != NULL_TREE)
+ {
+ while (DECL_CONTEXT (requires_decl) != NULL_TREE
+ && TREE_CODE (requires_decl) != TRANSLATION_UNIT_DECL)
+ requires_decl = DECL_CONTEXT (requires_decl);
+ if (requires_decl != NULL_TREE)
+ fn1 = IDENTIFIER_POINTER (DECL_NAME (requires_decl));
+ }
+
+ const char *fn2 = file_data->file_name;
+ if (tmp_decl != NULL_TREE)
+ {
+ while (DECL_CONTEXT (tmp_decl) != NULL_TREE
+ && TREE_CODE (tmp_decl) != TRANSLATION_UNIT_DECL)
+ tmp_decl = DECL_CONTEXT (tmp_decl);
+ if (tmp_decl != NULL_TREE)
+ fn2 = IDENTIFIER_POINTER (DECL_NAME (requires_decl));
+ }
+
+ char buf1[sizeof ("unified_address, unified_shared_memory, "
+ "reverse_offload")];
+ char buf2[sizeof ("unified_address, unified_shared_memory, "
+ "reverse_offload")];
+ omp_requires_to_name (buf2, sizeof (buf2),
+ val != OMP_REQUIRES_TARGET_USED
+ ? val
+ : (HOST_WIDE_INT) omp_requires_mask);
+ if (val != OMP_REQUIRES_TARGET_USED
+ && omp_requires_mask != OMP_REQUIRES_TARGET_USED)
+ {
+ omp_requires_to_name (buf1, sizeof (buf1),
+ omp_requires_mask);
+ error ("OpenMP %<requires%> directive with non-identical "
+ "clauses in multiple compilation units: %qs vs. "
+ "%qs", buf1, buf2);
+ inform (UNKNOWN_LOCATION, "%qs has %qs", fn1, buf1);
+ inform (UNKNOWN_LOCATION, "%qs has %qs", fn2, buf2);
+ }
+ else
+ {
+ error ("OpenMP %<requires%> directive with %qs specified "
+ "only in some compilation units", buf2);
+ inform (UNKNOWN_LOCATION, "%qs has %qs",
+ val != OMP_REQUIRES_TARGET_USED ? fn2 : fn1,
+ buf2);
+ inform (UNKNOWN_LOCATION, "but %qs has not",
+ val != OMP_REQUIRES_TARGET_USED ? fn1 : fn2);
+ }
+ error_emitted = true;
+ }
}
else
fatal_error (input_location,
@@ -1821,6 +1924,18 @@ input_offload_tables (bool do_force_output)
lto_destroy_simple_input_block (file_data, LTO_section_offload_table,
ib, data, len);
}
+#ifdef ACCEL_COMPILER
+ char *omp_requires_file = getenv ("GCC_OFFLOAD_OMP_REQUIRES_FILE");
+ if (omp_requires_file == NULL || omp_requires_file[0] == '\0')
+ fatal_error (input_location, "GCC_OFFLOAD_OMP_REQUIRES_FILE unset");
+ FILE *f = fopen (omp_requires_file, "wb");
+ if (!f)
+ fatal_error (input_location, "Cannot open omp_requires file %qs",
+ omp_requires_file);
+ uint32_t req_mask = omp_requires_mask;
+ fwrite (&req_mask, sizeof (req_mask), 1, f);
+ fclose (f);
+#endif
}
/* True when we need optimization summary for NODE. */
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 90a67810a89..8ea8cf4fd54 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -13573,6 +13573,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gcc_unreachable ();
}
+ /* Ensure that requires map is written via output_offload_tables, even if only
+ 'target (enter/exit) data' is used in the translation unit. */
+ if (ENABLE_OFFLOADING && (omp_requires_mask & OMP_REQUIRES_TARGET_USED))
+ g->have_offload = true;
+
clauses = gimple_omp_target_clauses (stmt);
gimple_seq dep_ilist = NULL;
diff --git a/gcc/omp-offload.cc b/gcc/omp-offload.cc
index 2b9aa87dbbf..ac236482db5 100644
--- a/gcc/omp-offload.cc
+++ b/gcc/omp-offload.cc
@@ -440,24 +440,6 @@ 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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 4ff58518d39..13141b9768e 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,25 @@
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-07-04 Tobias Burnus <tobias@codesourcery.com>
+ Chung-Lin Tang <cltang@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-c++-common/gomp/requires-4.c: Update dg-*.
+ * c-c++-common/gomp/reverse-offload-1.c: Likewise.
+ * c-c++-common/gomp/target-device-ancestor-2.c: Likewise.
+ * c-c++-common/gomp/target-device-ancestor-3.c: Likewise.
+ * c-c++-common/gomp/target-device-ancestor-4.c: Likewise.
+ * c-c++-common/gomp/target-device-ancestor-5.c: Likewise.
+ * gfortran.dg/gomp/target-device-ancestor-3.f90: Likewise.
+ * gfortran.dg/gomp/target-device-ancestor-4.f90: Likewise.
+ * gfortran.dg/gomp/target-device-ancestor-5.f90: Likewise.
+ * gfortran.dg/gomp/target-device-ancestor-2.f90: Likewise. Move
+ post-FE checks to ...
+ * gfortran.dg/gomp/target-device-ancestor-2a.f90: ... this new file.
+ * gfortran.dg/gomp/requires-8.f90: Update as we don't regard
+ 'declare target' for the 'requires' usage requirement.
+
2022-07-05 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline:
diff --git a/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c b/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c
index 9a3fa5230f8..3452156f948 100644
--- a/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c
@@ -43,7 +43,7 @@ tg_fn (int *x, int *y)
x2 = x2 + 2 + called_in_target1 ();
y2 = y2 + 7;
- #pragma omp target device(ancestor : 1) map(tofrom: x2)
+ #pragma omp target device(ancestor : 1) map(tofrom: x2) /* { dg-message "sorry, unimplemented: 'ancestor' not yet supported" } */
check_offload(&x2, &y2);
if (x2 != 2+2+3+42 || y2 != 3 + 7)
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c
index 6e9f02d8006..b16e701bd5a 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c
@@ -5,9 +5,7 @@
void
foo (int n)
{
- /* The following test is marked with 'xfail' because a previous 'sorry' from
- 'reverse_offload' suppresses the 'sorry' for 'ancestor'. */
- #pragma omp target device (ancestor: 1) /* { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } } */
+ #pragma omp target device (ancestor: 1)
;
@@ -19,9 +17,9 @@ foo (int n)
#pragma omp target device (ancestor : 42) /* { dg-error "the 'device' clause expression must evaluate to '1'" } */
;
- #pragma omp target device (ancestor : n) /* { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } } */
+ #pragma omp target device (ancestor : n)
;
- #pragma omp target device (ancestor : n + 1) /* { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } } */
+ #pragma omp target device (ancestor : n + 1)
;
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c
index b6ff84bcdab..b1520ff0636 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c
@@ -1,4 +1,4 @@
-#pragma omp requires reverse_offload /* { dg-message "sorry, unimplemented: 'reverse_offload' clause on 'requires' directive not supported yet" } */
+#pragma omp requires reverse_offload
void
foo ()
diff --git a/gcc/testsuite/gfortran.dg/gomp/requires-8.f90 b/gcc/testsuite/gfortran.dg/gomp/requires-8.f90
index e84d609ad29..583c5a56b32 100644
--- a/gcc/testsuite/gfortran.dg/gomp/requires-8.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/requires-8.f90
@@ -1,3 +1,7 @@
+module m0
+ integer :: x
+end module m0
+
module m ! { dg-error "has OpenMP device constructs/routines but does not set !.OMP REQUIRES UNIFIED_SHARED_MEMORY but other program units do" }
!$omp requires reverse_offload
contains
@@ -13,10 +17,14 @@ contains
end subroutine foo
end module m
-subroutine bar ! { dg-error "has OpenMP device constructs/routines but does not set !.OMP REQUIRES REVERSE_OFFLOAD but other program units do" }
+subroutine bar
!use m
- !$omp requires unified_shared_memory
+ !$omp requires unified_shared_memory ! Possibly OK - needs OpenMP Lang Spec clarification (-> #3240)
!$omp declare target
end subroutine bar
-! { dg-prune-output "not yet supported" }
+subroutine foobar ! { dg-error "has OpenMP device constructs/routines but does not set !.OMP REQUIRES REVERSE_OFFLOAD but other program units do" }
+ use m0
+ !$omp requires unified_shared_memory
+ !$omp target enter data map(to:x)
+end subroutine foobar
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2a.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2a.f90
new file mode 100644
index 00000000000..feb76fe2144
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2a.f90
@@ -0,0 +1,80 @@
+! { dg-do compile }
+
+implicit none
+
+integer :: a, b, c
+
+!$omp requires reverse_offload
+
+!$omp target device (ancestor: 1)
+!$omp end target
+
+!$omp target device (ancestor : a)
+!$omp end target
+
+!$omp target device (ancestor : a + 1)
+!$omp end target
+
+
+!$omp target device (device_num:42)
+!$omp end target
+
+!$omp target device (42)
+!$omp end target
+
+
+! Ensure that no OpenMP constructs appear inside target regions with 'ancestor'.
+
+!$omp target device (ancestor: 1)
+ !$omp teams ! { dg-error "OpenMP constructs are not allowed in target region with 'ancestor'" }
+ !$omp end teams
+!$omp end target
+
+!$omp target device (device_num: 1)
+ !$omp teams
+ !$omp end teams
+!$omp end target
+
+!$omp target device (1)
+ !$omp teams
+ !$omp end teams
+!$omp end target
+
+
+! Ensure that with 'ancestor' only the 'device', 'firstprivate', 'private',
+! 'defaultmap', and 'map' clauses appear on the construct.
+
+!$omp target nowait device (ancestor: 1) ! { dg-error "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" }
+!$omp end target
+
+!$omp target device (ancestor: 1) nowait ! { dg-error "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" }
+!$omp end target
+
+!$omp target nowait device (device_num: 1)
+!$omp end target
+
+!$omp target nowait device (1)
+!$omp end target
+
+!$omp target device (ancestor: 1) firstprivate (a) private (b) defaultmap (none) map (c)
+!$omp end target
+
+
+! Ensure that 'ancestor' is only used with 'target' constructs (not with
+! 'target data', 'target update' etc.).
+! The following test case is marked with 'xfail' because a previous 'sorry' from
+! 'reverse_offload' suppresses the 'sorry' for 'ancestor'.
+
+!$omp target data map (a) device (ancestor: 1) ! { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" }
+!$omp end target data
+
+!$omp target enter data map (to: a) device (ancestor: 1) ! { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" }
+!$omp target exit data map (from: a) device (ancestor: 1) ! { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" }
+
+!$omp target update to (a) device (ancestor: 1) ! { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" }
+
+!$omp target device (ancestor: 1) if(.false.)
+! { dg-error "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" "" { target *-*-* } .-1 }
+!$omp end target
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-5.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-5.f90
index 06a11eb5092..ca8d4b282a0 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-5.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-5.f90
@@ -6,7 +6,7 @@
!
module m
- !$omp requires reverse_offload ! { dg-error "REQUIRES directive is not yet supported" }
+ !$omp requires reverse_offload
contains
subroutine foo()
!$omp target device(ancestor:1)
@@ -17,7 +17,7 @@ contains
block
block
block
- !$omp target device(ancestor:1)
+ !$omp target device(ancestor:1) ! { dg-message "sorry, unimplemented: 'ancestor' not yet supported" }
!$omp end target
end block
end block
@@ -26,7 +26,7 @@ contains
end module m
subroutine foo()
- !$omp requires reverse_offload ! { dg-error "REQUIRES directive is not yet supported" }
+ !$omp requires reverse_offload
block
block
block
@@ -49,7 +49,7 @@ contains
end subroutine foo
program main
- !$omp requires reverse_offload ! { dg-error "REQUIRES directive is not yet supported" }
+ !$omp requires reverse_offload
contains
subroutine foo()
!$omp target device(ancestor:1)
diff --git a/include/ChangeLog.omp b/include/ChangeLog.omp
index aaf8892bd52..8b361c0d70b 100644
--- a/include/ChangeLog.omp
+++ b/include/ChangeLog.omp
@@ -1,3 +1,13 @@
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-07-04 Tobias Burnus <tobias@codesourcery.com>
+ Chung-Lin Tang <cltang@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * gomp-constants.h (GOMP_VERSION): Bump to 2.
+ (GOMP_REQUIRES_TARGET_USED): New defines.
+
2022-07-05 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline:
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 3c921f150f8..32b734bdbfc 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -309,7 +309,7 @@ enum gomp_map_kind
/* Versions of libgomp and device-specific plugins. GOMP_VERSION
should be incremented whenever an ABI-incompatible change is introduced
to the plugin interface defined in libgomp/libgomp.h. */
-#define GOMP_VERSION 1
+#define GOMP_VERSION 2
#define GOMP_VERSION_NVIDIA_PTX 1
#define GOMP_VERSION_INTEL_MIC 0
#define GOMP_VERSION_GCN 2
@@ -373,6 +373,7 @@ enum gomp_map_kind
#define GOMP_REQUIRES_UNIFIED_ADDRESS 0x10
#define GOMP_REQUIRES_UNIFIED_SHARED_MEMORY 0x20
#define GOMP_REQUIRES_REVERSE_OFFLOAD 0x80
+#define GOMP_REQUIRES_TARGET_USED 0x200
/* HSA specific data structures. */
diff --git a/libgcc/ChangeLog.omp b/libgcc/ChangeLog.omp
index 3fb6391410e..d8dfecbf796 100644
--- a/libgcc/ChangeLog.omp
+++ b/libgcc/ChangeLog.omp
@@ -1,3 +1,13 @@
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Reverted:
+ 2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * 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.
+
2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
* offloadstuff.c (__requires_mask_table): New symbol to mark start of
diff --git a/libgcc/offloadstuff.c b/libgcc/offloadstuff.c
index b2282924fb4..10e1fe19c8e 100644
--- a/libgcc/offloadstuff.c
+++ b/libgcc/offloadstuff.c
@@ -54,9 +54,6 @@ 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]
@@ -66,9 +63,6 @@ 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[];
@@ -83,9 +77,6 @@ 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/ChangeLog.omp b/libgomp/ChangeLog.omp
index b4d0d2bbb90..506f8f55944 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,68 @@
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-07-04 Tobias Burnus <tobias@codesourcery.com>
+ Chung-Lin Tang <cltang@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * libgomp-plugin.h (GOMP_OFFLOAD_get_num_devices): Add
+ omp_requires_mask arg.
+ * plugin/plugin-gcn.c (GOMP_OFFLOAD_get_num_devices): Likewise;
+ return -1 when device available but omp_requires_mask != 0.
+ * plugin/plugin-nvptx.c (GOMP_OFFLOAD_get_num_devices): Likewise.
+ * oacc-host.c (host_get_num_devices, host_openacc_get_property):
+ Update call.
+ * oacc-init.c (resolve_device, acc_init_1, acc_shutdown_1,
+ goacc_attach_host_thread_to_device, acc_get_num_devices,
+ acc_set_device_num, get_property_any): Likewise.
+ * target.c (omp_requires_mask): New global var.
+ (gomp_requires_to_name): New.
+ (GOMP_offload_register_ver): Handle passed omp_requires_mask.
+ (gomp_target_init): Handle omp_requires_mask.
+ * libgomp.texi (OpenMP 5.0): Update requires impl. status.
+ (OpenMP 5.1): Add a missed item.
+ (OpenMP 5.2): Mark linear-clause change as supported in C/C++.
+ * testsuite/libgomp.c-c++-common/requires-1-aux.c: New test.
+ * testsuite/libgomp.c-c++-common/requires-1.c: New test.
+ * testsuite/libgomp.c-c++-common/requires-2-aux.c: New test.
+ * testsuite/libgomp.c-c++-common/requires-2.c: New test.
+ * testsuite/libgomp.c-c++-common/requires-3-aux.c: New test.
+ * testsuite/libgomp.c-c++-common/requires-3.c: New test.
+ * testsuite/libgomp.c-c++-common/requires-4-aux.c: New test.
+ * testsuite/libgomp.c-c++-common/requires-4.c: New test.
+ * testsuite/libgomp.c-c++-common/requires-5-aux.c: New test.
+ * testsuite/libgomp.c-c++-common/requires-5.c: New test.
+ * testsuite/libgomp.c-c++-common/requires-6.c: New test.
+ * testsuite/libgomp.c-c++-common/requires-7-aux.c: New test.
+ * testsuite/libgomp.c-c++-common/requires-7.c: New test.
+ * testsuite/libgomp.fortran/requires-1-aux.f90: New test.
+ * testsuite/libgomp.fortran/requires-1.f90: New test.
+
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Reverted:
+ 2021-03-25 Thomas Schwinge <thomas@codesourcery.com>
+
+ * target.c (__requires_mask_table, __requires_mask_table_end): Add
+ '__attribute__((weak))'.
+
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Reverted:
+ 2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * 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.
+
2022-07-05 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline:
diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h
index 441ad172c01..e54f94cbcb5 100644
--- a/libgomp/libgomp-plugin.h
+++ b/libgomp/libgomp-plugin.h
@@ -125,10 +125,9 @@ extern void GOMP_PLUGIN_fatal (const char *, ...)
extern const char *GOMP_OFFLOAD_get_name (void);
extern unsigned int GOMP_OFFLOAD_get_caps (void);
extern int GOMP_OFFLOAD_get_type (void);
-extern int GOMP_OFFLOAD_get_num_devices (void);
+extern int GOMP_OFFLOAD_get_num_devices (unsigned int);
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 77669ed95b0..8d8fb3f3190 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1237,7 +1237,6 @@ 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/libgomp.texi b/libgomp/libgomp.texi
index df215fba4fa..527d457d8a9 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -189,8 +189,8 @@ The OpenMP 4.5 specification is fully supported.
env variable @tab Y @tab
@item Nested-parallel changes to @emph{max-active-levels-var} ICV @tab Y @tab
@item @code{requires} directive @tab P
- @tab Only fulfillable requirement are @code{atomic_default_mem_order}
- and @code{dynamic_allocators}
+ @tab complete but no non-host devices provides @code{unified_address},
+ @code{unified_shared_memory} or @code{reverse_offload}
@item @code{teams} construct outside an enclosing target region @tab Y @tab
@item Non-rectangular loop nests @tab Y @tab
@item @code{!=} as relational-op in canonical loop form for C/C++ @tab Y @tab
@@ -344,6 +344,8 @@ The OpenMP 4.5 specification is fully supported.
@item @code{unconstrained} and @code{reproducible} modifiers on @code{order}
clause @tab Y @tab
@item Support @code{begin/end declare target} syntax in C/C++ @tab N @tab
+@item Pointer predetermined firstprivate getting initialized
+to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab
@end multitable
@@ -353,6 +355,7 @@ The OpenMP 4.5 specification is fully supported.
@unnumberedsubsec New features listed in Appendix B of the OpenMP specification
@multitable @columnfractions .60 .10 .25
+@headitem Description @tab Status @tab Comments
@item @code{omp_in_explicit_task} routine and @emph{implicit-task-var} ICV
@tab N @tab
@item @code{omp}/@code{ompx}/@code{omx} sentinels and @code{omp_}/@code{ompx_}
diff --git a/libgomp/oacc-host.c b/libgomp/oacc-host.c
index 0ac70af8111..e3ef5f2efc6 100644
--- a/libgomp/oacc-host.c
+++ b/libgomp/oacc-host.c
@@ -54,7 +54,7 @@ host_get_type (void)
}
static int
-host_get_num_devices (void)
+host_get_num_devices (unsigned int omp_requires_mask __attribute__((unused)))
{
return 1;
}
@@ -71,12 +71,6 @@ 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)
{
@@ -244,7 +238,7 @@ host_openacc_get_property (int n, enum goacc_property prop)
{
union goacc_property_value nullval = { .val = 0 };
- if (n >= host_get_num_devices ())
+ if (n >= host_get_num_devices (0))
return nullval;
switch (prop)
@@ -288,7 +282,6 @@ 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/oacc-init.c b/libgomp/oacc-init.c
index 18d682ebd41..b6d1bc35d96 100644
--- a/libgomp/oacc-init.c
+++ b/libgomp/oacc-init.c
@@ -148,7 +148,7 @@ resolve_device (acc_device_t d, bool fail_is_error)
if (dispatchers[d]
&& !strcasecmp (goacc_device_type,
get_openacc_name (dispatchers[d]->name))
- && dispatchers[d]->get_num_devices_func () > 0)
+ && dispatchers[d]->get_num_devices_func (0) > 0)
goto found;
if (fail_is_error)
@@ -169,7 +169,7 @@ resolve_device (acc_device_t d, bool fail_is_error)
case acc_device_not_host:
/* Find the first available device after acc_device_not_host. */
while (known_device_type_p (++d))
- if (dispatchers[d] && dispatchers[d]->get_num_devices_func () > 0)
+ if (dispatchers[d] && dispatchers[d]->get_num_devices_func (0) > 0)
goto found;
if (d_arg == acc_device_default)
{
@@ -302,7 +302,7 @@ acc_init_1 (acc_device_t d, acc_construct_t parent_construct, int implicit)
base_dev = resolve_device (d, true);
- ndevs = base_dev->get_num_devices_func ();
+ ndevs = base_dev->get_num_devices_func (0);
if (ndevs <= 0 || goacc_device_num >= ndevs)
acc_dev_num_out_of_range (d, goacc_device_num, ndevs);
@@ -351,7 +351,7 @@ acc_shutdown_1 (acc_device_t d)
/* Get the base device for this device type. */
base_dev = resolve_device (d, true);
- ndevs = base_dev->get_num_devices_func ();
+ ndevs = base_dev->get_num_devices_func (0);
/* Unload all the devices of this type that have been opened. */
for (i = 0; i < ndevs; i++)
@@ -520,7 +520,7 @@ goacc_attach_host_thread_to_device (int ord)
base_dev = cached_base_dev;
}
- num_devices = base_dev->get_num_devices_func ();
+ num_devices = base_dev->get_num_devices_func (0);
if (num_devices <= 0 || ord >= num_devices)
acc_dev_num_out_of_range (acc_device_type (base_dev->type), ord,
num_devices);
@@ -599,7 +599,7 @@ acc_get_num_devices (acc_device_t d)
if (!acc_dev)
return 0;
- n = acc_dev->get_num_devices_func ();
+ n = acc_dev->get_num_devices_func (0);
if (n < 0)
n = 0;
@@ -779,7 +779,7 @@ acc_set_device_num (int ord, acc_device_t d)
cached_base_dev = base_dev = resolve_device (d, true);
- num_devices = base_dev->get_num_devices_func ();
+ num_devices = base_dev->get_num_devices_func (0);
if (num_devices <= 0 || ord >= num_devices)
acc_dev_num_out_of_range (d, ord, num_devices);
@@ -824,7 +824,7 @@ get_property_any (int ord, acc_device_t d, acc_device_property_t prop)
struct gomp_device_descr *dev = resolve_device (d, true);
- int num_devices = dev->get_num_devices_func ();
+ int num_devices = dev->get_num_devices_func (0);
if (num_devices <= 0 || ord >= num_devices)
acc_dev_num_out_of_range (d, ord, num_devices);
diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c
index ac40d6e7e2e..78956297d64 100644
--- a/libgomp/plugin/plugin-gcn.c
+++ b/libgomp/plugin/plugin-gcn.c
@@ -3221,10 +3221,16 @@ GOMP_OFFLOAD_version (void)
/* Return the number of GCN devices on the system. */
int
-GOMP_OFFLOAD_get_num_devices (void)
+GOMP_OFFLOAD_get_num_devices (unsigned int omp_requires_mask)
{
if (!init_hsa_context ())
return 0;
+ /* Return -1 if no omp_requires_mask cannot be fulfilled but
+ devices were present. */
+ if (hsa_context.agent_count > 0
+ && (omp_requires_mask & ~(GOMP_REQUIRES_UNIFIED_ADDRESS
+ | GOMP_REQUIRES_UNIFIED_SHARED_MEMORY)))
+ return -1;
return hsa_context.agent_count;
}
@@ -3892,17 +3898,6 @@ GOMP_OFFLOAD_is_usm_ptr (void *ptr)
return usm_splay_tree_lookup (&usm_map, &key);
}
-/* 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);
-}
-
/* }}} */
/* {{{ OpenACC Plugin API */
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 3f555b647f0..4ab88f465a0 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -1191,9 +1191,16 @@ GOMP_OFFLOAD_get_type (void)
}
int
-GOMP_OFFLOAD_get_num_devices (void)
-{
- return nvptx_get_num_devices ();
+GOMP_OFFLOAD_get_num_devices (unsigned int omp_requires_mask)
+{
+ int num_devices = nvptx_get_num_devices ();
+ /* Return -1 if no omp_requires_mask cannot be fulfilled but
+ devices were present. */
+ if (num_devices > 0
+ && (omp_requires_mask & ~(GOMP_REQUIRES_UNIFIED_ADDRESS
+ | GOMP_REQUIRES_UNIFIED_SHARED_MEMORY)))
+ return -1;
+ return num_devices;
}
bool
@@ -1264,17 +1271,6 @@ GOMP_OFFLOAD_fini_device (int n)
return true;
}
-/* 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);
-}
-
/* 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 df01958c324..0c5f83b9273 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -31,12 +31,12 @@
#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. */
#endif
#include <string.h>
+#include <stdio.h> /* For snprintf. */
#include <assert.h>
#include <errno.h>
@@ -99,17 +99,8 @@ 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. */
-__attribute__((weak))
-extern const unsigned int __requires_mask_table[];
-__attribute__((weak))
-extern const unsigned int __requires_mask_table_end[];
+/* OpenMP requires mask. */
+static int omp_requires_mask;
/* Similar to gomp_realloc, but release register_lock before gomp_fatal. */
@@ -2517,6 +2508,20 @@ gomp_unload_image_from_device (struct gomp_device_descr *devicep,
}
}
+static void
+gomp_requires_to_name (char *buf, size_t size, int requires_mask)
+{
+ char *end = buf + size, *p = buf;
+ if (requires_mask & GOMP_REQUIRES_UNIFIED_ADDRESS)
+ p += snprintf (p, end - p, "unified_address");
+ if (requires_mask & GOMP_REQUIRES_UNIFIED_SHARED_MEMORY)
+ p += snprintf (p, end - p, "%sunified_shared_memory",
+ (p == buf ? "" : ", "));
+ if (requires_mask & GOMP_REQUIRES_REVERSE_OFFLOAD)
+ p += snprintf (p, end - p, "%sreverse_offload",
+ (p == buf ? "" : ", "));
+}
+
/* This function should be called from every offload image while loading.
It gets the descriptor of the host func and var tables HOST_TABLE, TYPE of
the target, and TARGET_DATA needed by target plugin. */
@@ -2526,13 +2531,43 @@ GOMP_offload_register_ver (unsigned version, const void *host_table,
int target_type, const void *target_data)
{
int i;
+ int omp_req = 0;
if (GOMP_VERSION_LIB (version) > GOMP_VERSION)
gomp_fatal ("Library too old for offload (version %u < %u)",
GOMP_VERSION, GOMP_VERSION_LIB (version));
-
+
+ if (GOMP_VERSION_LIB (version) > 1)
+ {
+ omp_req = (int) (size_t) ((void **) target_data)[0];
+ target_data = &((void **) target_data)[1];
+ }
+
gomp_mutex_lock (®ister_lock);
+ if (omp_req && omp_requires_mask && omp_requires_mask != omp_req)
+ {
+ char buf1[sizeof ("unified_address, unified_shared_memory, "
+ "reverse_offload")];
+ char buf2[sizeof ("unified_address, unified_shared_memory, "
+ "reverse_offload")];
+ gomp_requires_to_name (buf2, sizeof (buf2),
+ omp_req != GOMP_REQUIRES_TARGET_USED
+ ? omp_req : omp_requires_mask);
+ if (omp_req != GOMP_REQUIRES_TARGET_USED
+ && omp_requires_mask != GOMP_REQUIRES_TARGET_USED)
+ {
+ gomp_requires_to_name (buf1, sizeof (buf1), omp_requires_mask);
+ gomp_fatal ("OpenMP 'requires' directive with non-identical clauses "
+ "in multiple compilation units: '%s' vs. '%s'",
+ buf1, buf2);
+ }
+ else
+ gomp_fatal ("OpenMP 'requires' directive with '%s' specified only in "
+ "some compilation units", buf2);
+ }
+ omp_requires_mask = omp_req;
+
/* Load image to all initialized devices. */
for (i = 0; i < num_devices; i++)
{
@@ -2621,20 +2656,6 @@ 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++)
{
@@ -4298,7 +4319,6 @@ 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);
@@ -4413,28 +4433,6 @@ 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
@@ -4461,8 +4459,30 @@ gomp_target_init (void)
if (gomp_load_plugin_for_device (¤t_device, plugin_name))
{
- new_num_devs = current_device.get_num_devices_func ();
- if (new_num_devs >= 1)
+ int omp_req = omp_requires_mask & ~GOMP_REQUIRES_TARGET_USED;
+ new_num_devs = current_device.get_num_devices_func (omp_req);
+ if (gomp_debug_var > 0 && new_num_devs < 0)
+ {
+ bool found = false;
+ int type = current_device.get_type_func ();
+ for (int img = 0; img < num_offload_images; img++)
+ if (type == offload_images[img].type)
+ found = true;
+ if (found)
+ {
+ char buf[sizeof ("unified_address, unified_shared_memory, "
+ "reverse_offload")];
+ gomp_requires_to_name (buf, sizeof (buf), omp_req);
+ char *name = (char *) malloc (cur_len + 1);
+ memcpy (name, cur, cur_len);
+ name[cur_len] = '\0';
+ gomp_debug (1,
+ "%s devices present but 'omp requires %s' "
+ "cannot be fulfilled", name, buf);
+ free (name);
+ }
+ }
+ else if (new_num_devs >= 1)
{
/* Augment DEVICES and NUM_DEVICES. */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-1-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-1-aux.c
index 8b9341523c6..bdca662e42f 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-1-aux.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-1-aux.c
@@ -1,6 +1,6 @@
/* { dg-skip-if "" { *-*-* } } */
-#pragma omp requires reverse_offload
+#pragma omp requires unified_address
int x;
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-1.c b/libgomp/testsuite/libgomp.c-c++-common/requires-1.c
index 0dd40bc0f59..fedf9779769 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-1.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-1.c
@@ -1,5 +1,8 @@
+/* { dg-do link { target { offload_target_nvptx || offload_target_amdgcn } } } */
/* { dg-additional-sources requires-1-aux.c } */
-/* { dg-require-effective-target omp_usm } */
+
+/* Check diagnostic by device-compiler's lto1.
+ Other file uses: 'requires unified_address'. */
#pragma omp requires unified_shared_memory
@@ -17,5 +20,5 @@ main (void)
return 0;
}
-/* { dg-output "libgomp: requires-directive clause inconsistency between compilation units detected" } */
-/* { dg-prune-output "device does not support required features" } */
+/* { dg-error "OpenMP 'requires' directive with non-identical clauses in multiple compilation units: 'unified_shared_memory' vs. 'unified_address'" "" { target *-*-* } 0 } */
+/* { dg-excess-errors "Ignore messages like: errors during merging of translation units|mkoffload returned 1 exit status" } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-2-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-2-aux.c
index 8b9341523c6..617577448ed 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-2-aux.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-2-aux.c
@@ -1,7 +1,5 @@
/* { dg-skip-if "" { *-*-* } } */
-#pragma omp requires reverse_offload
-
int x;
void foo (void)
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-2.c b/libgomp/testsuite/libgomp.c-c++-common/requires-2.c
index 0302f891c9c..be1830d0c46 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-2.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-2.c
@@ -1,6 +1,11 @@
+/* { dg-do link { target offloading_enabled } } */
+/* { dg-additional-options "-foffload=disable -flto" } */
/* { dg-additional-sources requires-2-aux.c } */
-#pragma omp requires reverse_offload
+/* Check diagnostic by host's lto1.
+ Other file does not have any 'omp requires'. */
+
+#pragma omp requires unified_shared_memory
int a[10];
extern void foo (void);
@@ -16,4 +21,5 @@ main (void)
return 0;
}
-/* { dg-output "libgomp: device does not support required features: reverse_offload" } */
+/* { dg-error "OpenMP 'requires' directive with 'unified_shared_memory' specified only in some compilation units" "" { target *-*-* } 0 } */
+/* { dg-excess-errors "Ignore messages like: errors during merging of translation units|mkoffload returned 1 exit status" } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-3-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-3-aux.c
new file mode 100644
index 00000000000..bdca662e42f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-3-aux.c
@@ -0,0 +1,11 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires unified_address
+
+int x;
+
+void foo (void)
+{
+ #pragma omp target
+ x = 1;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-3.c b/libgomp/testsuite/libgomp.c-c++-common/requires-3.c
new file mode 100644
index 00000000000..4b07ffdd09b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-3.c
@@ -0,0 +1,24 @@
+/* { dg-do link { target offloading_enabled } } */
+/* { dg-additional-sources requires-3-aux.c } */
+
+/* Check diagnostic by device-compiler's lto1.
+ Other file uses: 'requires unified_address'. */
+
+#pragma omp requires unified_address,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-error "OpenMP 'requires' directive with non-identical clauses in multiple compilation units: 'unified_address, unified_shared_memory' vs. 'unified_address'" "" { target *-*-* } 0 } */
+/* { dg-excess-errors "Ignore messages like: errors during merging of translation units|mkoffload returned 1 exit status" } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-4-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-4-aux.c
new file mode 100644
index 00000000000..b8b51ae8ca7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-4-aux.c
@@ -0,0 +1,13 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires reverse_offload
+
+/* Note: The file does not have neither of:
+ declare target directives, device constructs or device routines. */
+
+int x;
+
+void foo (void)
+{
+ x = 1;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-4.c b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
new file mode 100644
index 00000000000..128fdbb8463
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
@@ -0,0 +1,23 @@
+/* { dg-do link { target offloading_enabled } } */
+/* { dg-additional-options "-flto" } */
+/* { dg-additional-sources requires-4-aux.c } */
+
+/* Check diagnostic by device-compiler's or host compiler's lto1.
+ Other file uses: 'requires reverse_offload', but that's inactive as
+ there are no declare target directives, device constructs nor device routines */
+
+#pragma omp requires unified_address,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;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-5-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-5-aux.c
new file mode 100644
index 00000000000..d223749f0a1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-5-aux.c
@@ -0,0 +1,11 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires unified_shared_memory, unified_address, reverse_offload
+
+int x;
+
+void foo (void)
+{
+ #pragma omp target
+ x = 1;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-5.c b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
new file mode 100644
index 00000000000..c1e5540cfc5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
@@ -0,0 +1,21 @@
+/* { dg-do run { target { offload_target_nvptx || offload_target_amdgcn } } } */
+/* { dg-additional-sources requires-5-aux.c } */
+
+#pragma omp requires unified_shared_memory, unified_address, 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;
+}
+
+/* (Only) if GOMP_DEBUG=1, should print at runtime the following:
+ "devices present but 'omp requires unified_address, unified_shared_memory, reverse_offload' cannot be fulfilled" */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-6.c b/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
new file mode 100644
index 00000000000..b00c7459bbc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
@@ -0,0 +1,17 @@
+#pragma omp requires unified_shared_memory, unified_address, reverse_offload
+
+/* The requires line is not active as there is none of:
+ declare target directives, device constructs or device routines.
+ Thus, this code is expected to work everywhere. */
+
+int a[10];
+extern void foo (void);
+
+int
+main (void)
+{
+ for (int i = 0; i < 10; i++)
+ a[i] = 0;
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-7-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-7-aux.c
new file mode 100644
index 00000000000..0916db8a0ce
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-7-aux.c
@@ -0,0 +1,11 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires unified_address
+
+int x;
+
+void foo (void)
+{
+ x = 1;
+ #pragma omp target enter data map(always,to: x)
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-7.c b/libgomp/testsuite/libgomp.c-c++-common/requires-7.c
new file mode 100644
index 00000000000..c94a4c10846
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-7.c
@@ -0,0 +1,24 @@
+/* { dg-do link { target { offload_target_nvptx || offload_target_amdgcn } } } */
+/* { dg-additional-sources requires-7-aux.c } */
+
+/* Check diagnostic by device-compiler's lto1.
+ Other file uses: 'requires unified_address'. */
+
+#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-error "OpenMP 'requires' directive with non-identical clauses in multiple compilation units: 'unified_shared_memory' vs. 'unified_address'" "" { target *-*-* } 0 } */
+/* { dg-excess-errors "Ignore messages like: errors during merging of translation units|mkoffload returned 1 exit status" } */
diff --git a/libgomp/testsuite/libgomp.fortran/requires-1-aux.f90 b/libgomp/testsuite/libgomp.fortran/requires-1-aux.f90
new file mode 100644
index 00000000000..a18caeb4c69
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/requires-1-aux.f90
@@ -0,0 +1,14 @@
+! { dg-skip-if "" { *-*-* } }
+
+module m
+ integer x
+end module m
+
+subroutine foo
+ use m
+ implicit none
+ !$omp requires unified_address
+
+ x = 1
+ !$omp target enter data map(always,to: x)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/requires-1.f90 b/libgomp/testsuite/libgomp.fortran/requires-1.f90
new file mode 100644
index 00000000000..33741af15f1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/requires-1.f90
@@ -0,0 +1,26 @@
+! { dg-do link { target { offload_target_nvptx || offload_target_amdgcn } } }
+! { dg-additional-sources requires-1-aux.f90 }
+
+! Check diagnostic by device-compiler's lto1.
+! Other file uses: 'requires unified_address'.
+
+module m
+ integer :: a(10)
+ interface
+ subroutine foo
+ end
+ end interface
+end
+
+program main
+ !$omp requires unified_shared_memory
+
+ !$omp target
+ a = 0
+ !$omp end target
+
+ call foo ()
+end
+
+! { dg-error "OpenMP 'requires' directive with non-identical clauses in multiple compilation units: 'unified_shared_memory' vs. 'unified_address'" "" { target *-*-* } 0 }
+! { dg-excess-errors "Ignore messages like: errors during merging of translation units|mkoffload returned 1 exit status" }
diff --git a/liboffloadmic/ChangeLog.omp b/liboffloadmic/ChangeLog.omp
index e5d7b384670..e4ce22a3f00 100644
--- a/liboffloadmic/ChangeLog.omp
+++ b/liboffloadmic/ChangeLog.omp
@@ -1,3 +1,21 @@
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-07-04 Tobias Burnus <tobias@codesourcery.com>
+ Chung-Lin Tang <cltang@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_num_devices):
+ Return -1 when device available but omp_requires_mask != 0.
+
+2022-07-05 Tobias Burnus <tobias@codesourcery.com>
+
+ Reverted:
+ 2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_supported_features):
+ New function.
+
2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
* plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_supported_features):
diff --git a/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
index f92418fa416..33bae0650b4 100644
--- a/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
+++ b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
@@ -168,8 +168,12 @@ GOMP_OFFLOAD_get_type (void)
}
extern "C" int
-GOMP_OFFLOAD_get_num_devices (void)
+GOMP_OFFLOAD_get_num_devices (unsigned int omp_requires_mask)
{
+ /* Return -1 if no omp_requires_mask cannot be fulfilled but
+ devices were present. */
+ if (num_devices > 0 && omp_requires_mask != 0)
+ return -1;
TRACE ("(): return %d", num_devices);
return num_devices;
}
@@ -233,14 +237,6 @@ 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)
{
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-07-07 6:56 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-07-07 6:56 [gcc/devel/omp/gcc-12] OpenMP: Move omp requires checks to libgomp Tobias Burnus
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).