public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-12] nvptx/mkoffload.cc: Warn instead of error when reverse offload is not possible
@ 2022-09-12 13:50 Tobias Burnus
0 siblings, 0 replies; only message in thread
From: Tobias Burnus @ 2022-09-12 13:50 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:c8d2d18fc5d03590f6d7e452084b16cce5e044e4
commit c8d2d18fc5d03590f6d7e452084b16cce5e044e4
Author: Tobias Burnus <tobias@codesourcery.com>
Date: Mon Sep 12 15:47:20 2022 +0200
nvptx/mkoffload.cc: Warn instead of error when reverse offload is not possible
Reverse offload requests at least -misa=sm_35; with this patch, a warning
instead of an error is shown, still permitting reverse offload for all
other configured device types. This is achieved by not calling
GOMP_offload_register_ver (and stopping generating pointless 'static const char'
variables, once known.)
The tool_name as progname changes adds "nvptx " and "gcn " to the
"mkoffload: warning/error:" diagnostic.
gcc/ChangeLog:
* config/nvptx/mkoffload.cc (process): Replace a fatal_error by
a warning + not enabling offloading if -misa=sm_30 prevents
reverse offload.
(main): Use tool_name as progname for diagnostic.
* config/gcn/mkoffload.cc (main): Likewise.
libgomp/ChangeLog:
* libgomp.texi (Offload-Target Specifics: nvptx): Document
that reverse offload requires >= -march=sm_35.
* testsuite/libgomp.c-c++-common/requires-4.c: Build for nvptx
with -misa=sm_35.
* testsuite/libgomp.c-c++-common/requires-5.c: Likewise.
* testsuite/libgomp.c-c++-common/requires-6.c: Likewise.
* testsuite/libgomp.c-c++-common/reverse-offload-1.c: Likewise.
* testsuite/libgomp.fortran/reverse-offload-1.f90: Likewise.
* testsuite/libgomp.c/reverse-offload-sm30.c: New test.
(cherry picked from commit 6b43f556f392a7165582aca36a19fe7389d995b2)
Diff:
---
gcc/ChangeLog.omp | 11 +++++++++++
gcc/config/gcn/mkoffload.cc | 2 +-
gcc/config/nvptx/mkoffload.cc | 18 ++++++++++++++----
libgomp/ChangeLog.omp | 15 +++++++++++++++
libgomp/libgomp.texi | 3 +++
libgomp/testsuite/libgomp.c-c++-common/requires-4.c | 1 +
libgomp/testsuite/libgomp.c-c++-common/requires-5.c | 1 +
libgomp/testsuite/libgomp.c-c++-common/requires-6.c | 2 ++
.../testsuite/libgomp.c-c++-common/reverse-offload-1.c | 1 +
libgomp/testsuite/libgomp.c/reverse-offload-sm30.c | 15 +++++++++++++++
.../testsuite/libgomp.fortran/reverse-offload-1.f90 | 1 +
11 files changed, 65 insertions(+), 5 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index a2d1222bccf..4f80bcbd356 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,14 @@
+2022-09-12 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-09-12 Tobias Burnus <tobias@codesourcery.com>
+
+ * config/nvptx/mkoffload.cc (process): Replace a fatal_error by
+ a warning + not enabling offloading if -misa=sm_30 prevents
+ reverse offload.
+ (main): Use tool_name as progname for diagnostic.
+ * config/gcn/mkoffload.cc (main): Likewise.
+
2022-09-12 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline:
diff --git a/gcc/config/gcn/mkoffload.cc b/gcc/config/gcn/mkoffload.cc
index 2c7a1feafd2..14d7eedc790 100644
--- a/gcc/config/gcn/mkoffload.cc
+++ b/gcc/config/gcn/mkoffload.cc
@@ -844,7 +844,7 @@ main (int argc, char **argv)
FILE *cfile = stdout;
const char *outname = 0;
- progname = "mkoffload";
+ progname = tool_name;
diagnostic_initialize (global_dc, 0);
obstack_init (&files_to_cleanup);
diff --git a/gcc/config/nvptx/mkoffload.cc b/gcc/config/nvptx/mkoffload.cc
index 834b2059aac..854cd72f3c7 100644
--- a/gcc/config/nvptx/mkoffload.cc
+++ b/gcc/config/nvptx/mkoffload.cc
@@ -324,9 +324,19 @@ process (FILE *in, FILE *out, uint32_t omp_requires)
{
if (sm_ver && sm_ver[0] == '3' && sm_ver[1] == '0'
&& sm_ver[2] == '\n')
- fatal_error (input_location,
- "%<omp requires reverse_offload%> requires at least "
- "%<sm_35%> for %<-misa=%>");
+ {
+ warning_at (input_location, 0,
+ "%<omp requires reverse_offload%> requires at "
+ "least %<sm_35%> for "
+ "%<-foffload-options=nvptx-none=-march=%> - disabling"
+ " offload-code generation for this device type");
+ /* As now an empty file is compiled and there is no call to
+ GOMP_offload_register_ver, this device type is effectively
+ disabled. */
+ fflush (out);
+ ftruncate (fileno (out), 0);
+ return;
+ }
sm_ver2 = sm_ver;
version2 = version;
}
@@ -526,7 +536,7 @@ main (int argc, char **argv)
FILE *out = stdout;
const char *outname = 0;
- progname = "mkoffload";
+ progname = tool_name;
diagnostic_initialize (global_dc, 0);
if (atexit (mkoffload_cleanup) != 0)
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 44913b6e2ac..c9827ba53c2 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,18 @@
+2022-09-12 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline:
+ 2022-09-12 Tobias Burnus <tobias@codesourcery.com>
+
+ * libgomp.texi (Offload-Target Specifics: nvptx): Document
+ that reverse offload requires >= -march=sm_35.
+ * testsuite/libgomp.c-c++-common/requires-4.c: Build for nvptx
+ with -misa=sm_35.
+ * testsuite/libgomp.c-c++-common/requires-5.c: Likewise.
+ * testsuite/libgomp.c-c++-common/requires-6.c: Likewise.
+ * testsuite/libgomp.c-c++-common/reverse-offload-1.c: Likewise.
+ * testsuite/libgomp.fortran/reverse-offload-1.f90: Likewise.
+ * testsuite/libgomp.c/reverse-offload-sm30.c: New test.
+
2022-09-12 Tobias Burnus <tobias@codesourcery.com>
Backport from mainline:
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index c342aa6f5cc..0069b34f445 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -4381,6 +4381,9 @@ The implementation remark:
@item I/O within OpenMP target regions and OpenACC parallel/kernels is supported
using the C library @code{printf} functions and the Fortran
@code{print}/@code{write} statements.
+@item Compilation OpenMP code that contains @code{requires reverse_offload}
+ requires at least @code{-march=sm_35}, compiling for @code{-march=sm_30}
+ is not supported.
@end itemize
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-4.c b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
index 6ed5a5f647a..5883eff0d93 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
@@ -1,4 +1,5 @@
/* { dg-additional-options "-flto" } */
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
/* { dg-additional-sources requires-4-aux.c } */
/* Check no diagnostic by device-compiler's or host compiler's lto1.
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-5.c b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
index 7fe0c735d27..d43d78db6fa 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
/* { dg-additional-sources requires-5-aux.c } */
/* Depending on offload device capabilities, it may print something like the
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-6.c b/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
index b00c7459bbc..a25b4d2dedd 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
@@ -1,3 +1,5 @@
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
+
#pragma omp requires unified_shared_memory, unified_address, reverse_offload
/* The requires line is not active as there is none of:
diff --git a/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c b/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c
index 976e129f560..52d828caf1c 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c
@@ -1,4 +1,5 @@
/* { dg-do run } */
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
/* { dg-additional-sources reverse-offload-1-aux.c } */
/* Check that reverse offload works in particular:
diff --git a/libgomp/testsuite/libgomp.c/reverse-offload-sm30.c b/libgomp/testsuite/libgomp.c/reverse-offload-sm30.c
new file mode 100644
index 00000000000..fbfeae1fd41
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/reverse-offload-sm30.c
@@ -0,0 +1,15 @@
+/* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload-options=nvptx-none=-march=sm_30 -foffload=-mptx=_" } */
+
+#pragma omp requires reverse_offload
+
+int
+main ()
+{
+ #pragma omp target
+ {
+ }
+ return 0;
+}
+
+/* { dg-warning "'omp requires reverse_offload' requires at least 'sm_35' for '-foffload-options=nvptx-none=-march=' - disabling offload-code generation for this device type" "" { target *-*-* } 0 } */
diff --git a/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90 b/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90
index 7cfb8b6552e..de68011f8f7 100644
--- a/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90
@@ -1,4 +1,5 @@
! { dg-do run }
+! { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } }
! { dg-additional-sources reverse-offload-1-aux.f90 }
! Check that reverse offload works in particular:
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-09-12 13:50 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-09-12 13:50 [gcc/devel/omp/gcc-12] nvptx/mkoffload.cc: Warn instead of error when reverse offload is not possible 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).